diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000..b71fba4 --- /dev/null +++ b/.gitignore @@ -0,0 +1,23 @@ +# ambienti python / cache +c/mio_env/ +__pycache__/ +**/__pycache__/ +*.pyc + +# binari compilati (si rigenerano con make / coli build) +c/glm +c/olmoe +c/iobench +c/tok_test + +# oracoli tiny generati (make_glm_oracle.py) e dati benchmark scaricati +c/glm_tiny/ +c/glm_tiny_i2/ +c/glm_tiny_i4/ +c/glm_tiny_mix/ +c/bench/ + +# pesi modello / artefatti di run +*.safetensors +stats*.txt +*.log diff --git a/README.md b/README.md index 66b60c6..19f42e0 100644 --- a/README.md +++ b/README.md @@ -1 +1,82 @@ -# colibri \ No newline at end of file +# colibrì 🐦 + +**Tiny engine, immense model.** Run **GLM-5.2 (744B-parameter MoE)** on a consumer machine with ~25 GB of RAM — in pure C, with zero dependencies, by streaming experts from disk. + +``` +$ ./coli chat + 🐦 colibrì v1.0 — GLM-5.2 · 744B MoE · int4 · streaming CPU + ✓ pronto in 32s · residente 9.9 GB + › ciao! + ◆ Ciao! 😊 Come posso aiutarti oggi? +``` + +## The idea + +A 744B Mixture-of-Experts model activates only ~40B parameters per token — and only ~11 GB of those change from token to token (the routed experts). So: + +- the **dense part** (attention, shared experts, embeddings — ~17B params) stays **resident in RAM at int4** (~9.9 GB); +- the **21,504 routed experts** (75 MoE layers × 256 experts + the MTP head, ~19 MB each at int4) live **on disk** (~370 GB) and are **streamed on demand**, with a per-layer LRU cache, an optional pinned hot-store, and the OS page cache as a free L2. + +The engine is a single C file (`c/glm.c`, ~1,300 lines) plus small headers. No BLAS, no Python at runtime, no GPU. + +## What's implemented + +- **Faithful GLM-5.2 (`glm_moe_dsa`) forward** — validated token-exact against a `transformers` oracle (teacher-forcing 32/32, greedy 20/20 on a tiny-random model with the real architecture). +- **MLA attention** (q/kv-LoRA, interleaved partial RoPE) with **compressed KV-cache**: 576 floats/token instead of 32,768 (57× smaller — GLM-5.2 has 64 heads and no GQA). +- **DeepSeek-V3-style sigmoid router** (noaux_tc, routed_scaling_factor), shared expert, first-3-dense layers. +- **Native MTP speculative decoding** — GLM-5.2's own multi-token-prediction head (layer 78) drafts tokens that the main model verifies in one batched forward. Measured **2.00 tokens/forward (100% acceptance)** on structured text. Lossless: output identical to greedy. +- **Quantization kernels**: int8 / packed int4 / packed int2, per-row scales, AVX2, dequant-on-use. Packing validated bit-identical to the int8 container. +- **Batch-union MoE**: in prefill (and MTP verification), each unique expert of the batch is read once and applied to every position that routes to it. +- **Byte-level BPE tokenizer in C** (GPT-2-style with Unicode-property regex, 320k merges). +- **RAM safety**: the expert cache is auto-sized from `MemAvailable` at startup — an honest peak projection (working set, KV, MTP row, reconstruction buffers) so the kernel OOM-killer never fires. +- **Offline FP8→int4 converter** (`c/convert_fp8_to_int4.py`): downloads one shard at a time (~5 GB), dequants (128×128 block scales), requantizes to the engine's container, deletes the shard — the 756 GB FP8 checkpoint never needs to exist on disk at once. Resumable. + +## Honest numbers (WSL2, 12 cores, 25 GB RAM, NVMe via VHDX) + +| metric | value | +|---|---| +| model on disk (int4 container) | ~370 GB | +| resident RAM (dense, int4) | 9.9 GB | +| load time | ~30 s | +| peak RSS during chat | ~20 GB (auto-capped) | +| cold decode cost | ~11 GB disk reads/token (75 layers × 8 experts) | +| disk ceiling (VHDX random) | ~1 GB/s → ~0.05–0.1 tok/s cold | +| MTP speculation | 2.0 tok/forward measured | + +This is not fast. It is a 744B frontier-class model **answering correctly on a machine that costs less than one H100 fan**. Warm cache, pinned hot experts and MTP push the useful-response latency down considerably; the physics of the disk does the rest. + +## Quick start + +```bash +cd c +./setup.sh # checks gcc/OpenMP, builds, self-tests + +# convert the model (resumable, needs ~400 GB free on a real ext4/NVMe path): +./coli convert # from zai-org/GLM-5.2-FP8 + +# chat (RAM budget and expert cache size itself automatically): +COLI_MODEL=/path/to/glm52_i4 ./coli chat +``` + +Useful knobs (env or flags): `--topp 0.7` adaptive expert top-p (30–40% less disk), `--ngen N` max tokens, `STATS=f`/`PIN=f PIN_GB=g` record expert usage then pin the hottest in RAM, `THINK=1` enable GLM-5.2's reasoning block, `DRAFT=n` MTP draft depth, `TF=1` teacher-forcing validation. + +## Repo layout + +``` +c/glm.c the engine (GLM-5.2 forward, streaming MoE, MTP, serve mode) +c/st.h safetensors reader: pread + fadvise, no mmap (RSS stays flat) +c/tok.h byte-level BPE tokenizer in C +c/coli CLI: chat / run / bench / convert / info +c/convert_fp8_to_int4.py disk-safe FP8 → int4 converter +c/make_glm_oracle.py tiny-random oracle generator for validation +c/olmoe.c stage-A engine (OLMoE), first validation target +*.py research scripts (cost model, trace analysis, py engine) +``` + +## Why "colibrì" + +The hummingbird weighs a few grams, hovers in place, and visits a thousand flowers a day. This engine keeps a 744-billion-parameter giant alive on hummingbird rations: 25 GB of RAM, twelve CPU cores, and a lot of disk patience. + +## License + +Apache 2.0. GLM-5.2 weights are released by Z.ai under MIT. diff --git a/c/Makefile b/c/Makefile new file mode 100644 index 0000000..b035b9e --- /dev/null +++ b/c/Makefile @@ -0,0 +1,22 @@ +CC = gcc +# ARCH=native -> ottimizzato per QUESTA macchina (default, piu' veloce). +# ARCH=x86-64-v3 -> binario PORTABILE su qualsiasi x86-64 moderno con AVX2 (per distribuire). +# ARCH=x86-64 -> massima compatibilita' (niente AVX2: usa il path scalare di fallback). +ARCH ?= native +CFLAGS = -O3 -march=$(ARCH) -fopenmp -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function +LDFLAGS = -lm -fopenmp + +all: glm + +glm: glm.c st.h json.h tok.h tok_unicode.h + $(CC) $(CFLAGS) glm.c -o glm $(LDFLAGS) + +olmoe: olmoe.c st.h json.h + $(CC) $(CFLAGS) olmoe.c -o olmoe $(LDFLAGS) + +# binario portabile da distribuire su altre macchine x86-64 +portable: + $(MAKE) glm ARCH=x86-64-v3 + +clean: + rm -f olmoe glm diff --git a/c/coli b/c/coli new file mode 100644 index 0000000..bae58d0 --- /dev/null +++ b/c/coli @@ -0,0 +1,291 @@ +#!/usr/bin/env python3 +""" +colibrì — piccolo motore, modello immenso. +CLI per far girare GLM-5.2 (744B) in locale, su CPU, in ~15-26 GB di RAM. + + coli chat chat interattiva (carica il modello UNA volta) + coli run "prompt" generazione singola + coli info stato: modello, RAM, disco, config + coli bench [task...] benchmark di qualità (MMLU/HellaSwag/...) + coli convert converte GLM-5.2-FP8 -> int4 (streaming) + coli build compila il motore + +Config via env o flag (validi anche dopo il sottocomando): + COLI_MODEL= modello (default /home/vincenzo/glm52_i4) + --ram N budget RAM in GB (auto-cap cache expert) + --topp P top-p adattivo sugli expert --topk N top-k fisso + --ngen N token massimi per risposta --cap N slot cache/layer +""" +import os, sys, subprocess, argparse, json, time, signal, shutil, threading, re, codecs, tempfile + +HERE = os.path.dirname(os.path.abspath(__file__)) +GLM = os.path.join(HERE, "glm") +DEF_MODEL = os.environ.get("COLI_MODEL", "/home/vincenzo/glm52_i4") +END = b"\x01\x01END\x01\x01\n" +READY = b"\x01\x01READY\x01\x01\n" + +# ---------- palette & stile ---------- +def _c(n): return f"\033[38;5;{n}m" +class C: + teal=_c(37); cyan=_c(80); mag=_c(170); org=_c(208); grn=_c(78); yel=_c(179) + dim="\033[2m"; b="\033[1m"; r="\033[0m"; gray=_c(242); dgray=_c(238) + @staticmethod + def off(): + for k,v in vars(C).items(): + if isinstance(v,str) and v.startswith("\033"): setattr(C,k,"") +TTY = sys.stdout.isatty() or os.environ.get("COLI_COLOR")=="1" +if not TTY: C.off() + +# ---------- colibrì 8-bit (pixel art, 2 pixel verticali per carattere) ---------- +SPRITE = [ + "....MMM.........", + "...MMMMM..w.....", + "....MMMM.ww.....", + "OOOOTTeTCC......", + "....TTTTTCC.....", + ".....TTTTCC.....", + "......TTCC......", + ".......TC.......", + "........C.......", + "................", +] +PAL = {"M":170, "T":37, "C":80, "O":208, "e":231, "w":80, ".":None} + +def sprite_lines(): + if not TTY: + return [" (\\ ", " )·> ", " / \\ ", " ", " "] + out=[] + for y in range(0,len(SPRITE),2): + top, bot = SPRITE[y], SPRITE[y+1] if y+1len(sentinel): + out=pend[:-len(sentinel)]; pend=pend[-len(sentinel):] + on_bytes(out) + +# ---------- comandi ---------- +def cmd_build(a): + banner("build") + sys.exit(subprocess.call(["make","-C",HERE,"glm"])) + +def cmd_info(a): + banner("info") + cfgp=os.path.join(a.model,"config.json") + def row(k,v): print(f" {C.gray}{k:<10}{C.r} {v}") + if os.path.exists(cfgp): + c=json.load(open(cfgp)) + row("modello", a.model) + row("arch", f"hidden {c.get('hidden_size')} · {c.get('num_hidden_layers')} layer · " + f"{c.get('n_routed_experts')} expert/layer · top-{c.get('num_experts_per_tok')}") + sts=[x for x in os.listdir(a.model) if x.endswith('.safetensors')] + sz=sum(os.path.getsize(os.path.join(a.model,x)) for x in sts) + row("shard", f"{len(sts)} file · {sz/1e9:.0f} GB su disco") + else: + print(f" {C.yel}config.json non presente (conversione incompleta?){C.r}") + try: + mi=open('/proc/meminfo').read() + tot=int(re.search(r'MemTotal:\s+(\d+)',mi).group(1))/1e6 + av=int(re.search(r'MemAvailable:\s+(\d+)',mi).group(1))/1e6 + row("RAM", f"{tot:.0f} GB totali · {av:.1f} GB disponibili") + except Exception: pass + fs=os.statvfs(a.model if os.path.isdir(a.model) else HERE) + row("disco", f"{fs.f_bavail*fs.f_frsize/1e9:.0f} GB liberi") + row("motore", "pronto ✓" if os.path.exists(GLM) else "da compilare (coli build)") + knobs=[] + if a.ram: knobs.append(f"ram {a.ram}GB") + if a.topp: knobs.append(f"topp {a.topp}") + if a.topk: knobs.append(f"topk {a.topk}") + if knobs: row("tuning", " · ".join(knobs)) + print() + +def cmd_run(a): + need_model(a.model) + prompt=" ".join(a.prompt) if a.prompt else sys.exit('uso: coli run "il tuo prompt"') + banner("run") + # template ufficiale GLM-5.2: niente \n dopo i ruoli; = risposta diretta (nothink) + e=env_for(a); e["PROMPT"]=f"[gMASK]<|user|>{prompt}<|assistant|>" + sys.exit(subprocess.call([GLM, str(a.cap)], env=e)) + +def cmd_chat(a): + need_model(a.model) + banner(f"chat · {os.path.basename(a.model)} · ram {a.ram or '-'}GB · topp {a.topp or 'off'}") + errlog=tempfile.NamedTemporaryFile(mode="w+", suffix=".log", delete=False) + e=env_for(a); e["SERVE"]="1" + p=subprocess.Popen([GLM,str(a.cap)], env=e, stdin=subprocess.PIPE, + stdout=subprocess.PIPE, stderr=errlog, bufsize=0) + sp=Spinner("sveglio il gigante (744B)…"); sp.start() + st=stream_turn(p, READY, lambda b: None) + sp.stop() + if st is None: + errlog.seek(0); print(errlog.read()[-1500:]); sys.exit("il motore è uscito durante il load") + errlog.flush() + try: + elog=open(errlog.name).read() + mload=re.search(r"caricato in ([0-9.]+)s \| densa residente: ([0-9.]+) MB", elog) + extra=" · ".join(l.strip() for l in elog.splitlines() if l.startswith("[RAM_GB") or l.startswith("[PIN]")) + if mload: print(f" {C.grn}✓{C.r} pronto in {mload.group(1)}s {C.dim}· residente {float(mload.group(2))/1000:.1f} GB · RSS {st.get('rss','?')} GB{C.r}") + if extra: print(f" {C.dgray}{extra}{C.r}") + except Exception: pass + print(f" {C.dim}scrivi e premi invio · :reset memoria · :q esci{C.r}\n") + w=term_w()-4 + try: + while True: + if TTY: + print(f" {C.dgray}╭{'─'*w}╮{C.r}") + try: msg=input(f" {C.dgray}│{C.r} {C.teal}{C.b}›{C.r} ") + except EOFError: print(); break + print(f" {C.dgray}╰{'─'*w}╯{C.r}") + else: + try: msg=input() + except EOFError: break + msg=msg.strip() + if msg in (":q",":quit","exit"): break + if not msg: continue + if msg==":reset": + p.stdin.write(b"\x02RESET\n"); p.stdin.flush() + stream_turn(p, END, lambda b: None) + print(f" {C.dim}✦ memoria azzerata{C.r}\n"); continue + p.stdin.write((msg.replace("\n"," ")+"\n").encode()); p.stdin.flush() + print(f"\n {C.teal}◆ colibrì{C.r}") + dec=codecs.getincrementaldecoder("utf-8")("replace") + state={"first":True} + sp2=Spinner("pensa…"); sp2.start() + def echo(bs, _dec=dec, _st=state): + if _st["first"]: + sp2.stop(); _st["first"]=False + sys.stdout.write(" ") + s=_dec.decode(bs) + if s: sys.stdout.write(s.replace("\n","\n ")); sys.stdout.flush() + t0=time.time() + st=stream_turn(p, END, echo) + sp2.stop() + if st is None: print(f"\n {C.yel}[motore terminato]{C.r}"); break + 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}\n") + else: + print() + except KeyboardInterrupt: + print(f"\n {C.dim}interrotto{C.r}") + finally: + try: p.stdin.close(); p.terminate() + except Exception: pass + try: os.unlink(errlog.name) + except Exception: pass + print(f" {C.teal}ciao{C.r} {C.dim}— il colibrì torna al nido{C.r} 🐦\n") + +def cmd_bench(a): + need_model(a.model) + banner("bench") + cmd=[sys.executable, os.path.join(HERE,"eval_glm.py"), "--snap",a.model, + "--tasks", ",".join(a.tasks) if a.tasks else "hellaswag,arc_challenge,mmlu", + "--limit", str(a.limit), "--data", a.data] + if a.ram: cmd+=["--ram",str(a.ram)] + e=dict(os.environ) + if a.topp: e["TOPP"]=str(a.topp) + if a.topk: e["TOPK"]=str(a.topk) + sys.exit(subprocess.call(cmd, env=e)) + +def cmd_convert(a): + banner("convert") + cmd=[sys.executable, os.path.join(HERE,"convert_fp8_to_int4.py"), + "--repo", a.repo, "--outdir", a.model, "--ebits", str(a.ebits), "--io-bits", str(a.io_bits)] + if a.xbits: cmd+=["--xbits",str(a.xbits)] + print(f" {C.dim}{' '.join(cmd)}{C.r}") + sys.exit(subprocess.call(cmd)) + +def main(): + common=argparse.ArgumentParser(add_help=False) + common.add_argument("--model", default=DEF_MODEL); common.add_argument("--ram", type=int, default=0) # 0 = auto (il motore usa l'88% della RAM disponibile) + common.add_argument("--cap", type=int, default=8); common.add_argument("--ngen", type=int, default=256) + common.add_argument("--topp", type=float, default=0); common.add_argument("--topk", type=int, default=0) + ap=argparse.ArgumentParser(prog="coli", parents=[common], description="colibrì — GLM-5.2 in locale") + sub=ap.add_subparsers(dest="cmd") + sub.add_parser("build", parents=[common]); sub.add_parser("info", parents=[common]) + pr=sub.add_parser("run", parents=[common]); pr.add_argument("prompt", nargs="*") + sub.add_parser("chat", parents=[common]) + pb=sub.add_parser("bench", parents=[common]); pb.add_argument("tasks", nargs="*") + pb.add_argument("--limit",type=int,default=40); pb.add_argument("--data",default=os.path.join(HERE,"bench")) + pc=sub.add_parser("convert", parents=[common]); pc.add_argument("--repo",default="zai-org/GLM-5.2-FP8") + pc.add_argument("--ebits",type=int,default=4); pc.add_argument("--io-bits",type=int,default=8); pc.add_argument("--xbits",type=int,default=0) + a=ap.parse_args() + {"build":cmd_build,"info":cmd_info,"run":cmd_run,"chat":cmd_chat,"bench":cmd_bench, + "convert":cmd_convert}.get(a.cmd, lambda _:(banner(),print(__doc__)))(a) + +if __name__=="__main__": + signal.signal(signal.SIGINT, signal.default_int_handler) + main() diff --git a/c/convert_fp8_to_int4.py b/c/convert_fp8_to_int4.py new file mode 100644 index 0000000..0348f6e --- /dev/null +++ b/c/convert_fp8_to_int4.py @@ -0,0 +1,250 @@ +""" +Convertitore GLM-5.2-FP8 -> nostro container int4 (STADIO B). + +Strategia DISK-SAFE (richiesta dell'utente): scarica UNO shard (~5 GB), lo converte in +int4, lo CANCELLA, passa al prossimo. Il disco non si riempie mai: picco = 1 shard + l'output +int4 che cresce fino a ~372 GB. Controllo di spazio che si ferma se manca margine. + +Cosa fa per ogni tensore: + - pesi FP8 (e4m3) con `*.weight_scale_inv` -> dequant a blocchi 128x128 -> f32 + - pesi BF16 (norme/embed/lm_head/...) -> f32 + poi: + - attn/mlp/shared/expert/embed/lm_head -> QUANTIZZATO int4 (o int8) con la STESSA matematica + del motore C (np.rint = lrintf, stesse soglie, stesso packing dei nibble) -> token identici + - norme / router (mlp.gate.weight) / bias / e_score_correction_bias -> tenuti F32 + - indexer DSA / layer MTP (78) / shared_head / eh_proj / *norm dell'indexer -> SALTATI + +Output: una dir di safetensors leggibile dal motore C (per ogni peso quantizzato: `nome` U8 = +dati impacchettati, `nome.qs` F32 = scale per riga). + +USO: + # test locale (oracolo tiny, niente download): converte una dir gia' presente + python3 convert_fp8_to_int4.py --indir glm_tiny --outdir glm_tiny_i4 --ebits 4 --io-bits 4 + # selftest del dequant fp8 (richiede torch) + python3 convert_fp8_to_int4.py --selftest + # reale: scarica+converte+cancella shard per shard + python3 convert_fp8_to_int4.py --repo zai-org/GLM-5.2-FP8 --outdir /home/vincenzo/glm52_i4 +""" +import os, sys, glob, json, shutil, argparse +import numpy as np + +# ---------- quantizzazione: identica al C (glm.c) ---------- +def quant_int8(w, bits): # w: [O,I] f32 -> (qbytes U8 [O*I], scale f32 [O]) + qmax = (1 << (bits - 1)) - 1 + amax = np.abs(w).max(axis=1, keepdims=True) + s = np.maximum(amax / qmax, 1e-8) + q = np.clip(np.rint(w / s), -qmax - 1, qmax).astype(np.int8) + return q.reshape(-1).view(np.uint8).copy(), s[:, 0].astype(np.float32) + +def quant_int4(w, bits): # -> (qbytes U8 [O*ceil(I/2)], scale f32 [O]) + O, I = w.shape + qmax = (1 << (bits - 1)) - 1 + amax = np.abs(w).max(axis=1, keepdims=True) + s = np.maximum(amax / qmax, 1e-8) + q = np.clip(np.rint(w / s), -8, qmax).astype(np.int32) # nibble [-8,7] + rb = (I + 1) // 2 + out = np.zeros((O, rb), np.uint8) + v0 = (q[:, 0::2] + 8).astype(np.uint8) + out[:, :v0.shape[1]] = v0 + if I > 1: + v1 = (q[:, 1::2] + 8).astype(np.uint8) + out[:, :v1.shape[1]] |= (v1 << 4) + return out.reshape(-1), s[:, 0].astype(np.float32) + +def quant_int2(w, bits): # -> (qbytes U8 [O*ceil(I/4)], scale f32 [O]); 4/byte + O, I = w.shape + qmax = (1 << (bits - 1)) - 1 # bits=2 -> qmax=1, valori [-2,1] + amax = np.abs(w).max(axis=1, keepdims=True) + s = np.maximum(amax / qmax, 1e-8) + q = np.clip(np.rint(w / s), -2, qmax).astype(np.int32) + rb = (I + 3) // 4 + out = np.zeros((O, rb), np.uint8) + for k in range(4): # impacchetta 4 valori per byte (identico a pack_int2 in C) + vk = q[:, k::4] + out[:, :vk.shape[1]] |= ((vk + 2).astype(np.uint8) << (k * 2)) + return out.reshape(-1), s[:, 0].astype(np.float32) + +# ---------- classificazione dei tensori ---------- +def layer_idx(name): + p = name.split(".") + if len(p) > 2 and p[0] == "model" and p[1] == "layers": + try: return int(p[2]) + except ValueError: return -1 + return -1 + +def classify(name, n_layers, keep_mtp=False): + if name.endswith("_scale_inv"): return "consumed" # gestito col suo peso + li = layer_idx(name) + if keep_mtp: + if li != n_layers: return "skip" # solo il layer MTP + if "indexer" in name: return "skip" # il DSA indexer resta un no-op + else: + if li >= n_layers: return "skip" # layer MTP (78) + if any(k in name for k in ["indexer", "indexers_proj", "eh_proj", + "enorm", "hnorm", "shared_head"]): return "skip" + if name.endswith("e_score_correction_bias"): return "f32" + if name.endswith("mlp.gate.weight"): return "f32" # router (NON gate_proj) + 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) + return "f32" + +# ---------- dequant di un tensore (fp8+scale a blocchi / bf16 / f32) ---------- +def dequant(f, name): + import torch + sl = f.get_slice(name); dt = sl.get_dtype() + if dt in ("F8_E4M3", "float8_e4m3fn"): + w = f.get_tensor(name).to(torch.float32) + sc = f.get_tensor(name + "_scale_inv").to(torch.float32) # [ceil(O/128),ceil(I/128)] + O, I = w.shape + sc = sc.repeat_interleave(128, 0).repeat_interleave(128, 1)[:O, :I] + return (w * sc).numpy() + return f.get_tensor(name).to(torch.float32).numpy() + +def convert_shard(path, out_dict, n_layers, ebits, io_bits, xbits, keep_mtp=False): + from safetensors import safe_open + with safe_open(path, framework="pt") as f: + for name in f.keys(): + kind = classify(name, n_layers, keep_mtp) + if kind in ("skip", "consumed"): continue + w = dequant(f, name) + if kind == "f32": + out_dict[name] = w.astype(np.float32) + else: + bits = io_bits if kind == "io" else xbits if kind == "x" else ebits + if w.ndim != 2: # es. bias 1D non previsto come 'q' -> tienilo f32 + out_dict[name] = w.astype(np.float32); continue + q, s = (quant_int2(w, bits) if bits <= 2 else + quant_int4(w, bits) if bits <= 4 else quant_int8(w, bits)) + out_dict[name] = q + out_dict[name + ".qs"] = s + +def free_gb(p): return shutil.disk_usage(p).free / 1e9 + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--repo", default=None) + ap.add_argument("--indir", default=None) + ap.add_argument("--outdir", required=False) + ap.add_argument("--ebits", type=int, default=4) # bit residenti: attn/dense-mlp/shared + 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 + ap.add_argument("--n-layers", type=int, default=78) + ap.add_argument("--min-free-gb", type=float, default=20.0) + ap.add_argument("--selftest", action="store_true") + ap.add_argument("--mtp", action="store_true", + help="scarica/converte SOLO la testa MTP (model.layers..*) -> out-mtp-*.safetensors") + a = ap.parse_args() + if a.xbits is None: a.xbits = a.ebits + + if a.selftest: + import torch + w = (torch.randn(256, 256) * 0.3) + O, I = w.shape; bs = 128 + sc = torch.zeros(O // bs, I // bs) + for bi in range(O // bs): + for bj in range(I // bs): + blk = w[bi*bs:(bi+1)*bs, bj*bs:(bj+1)*bs] + sc[bi, bj] = blk.abs().max() / 448.0 + q = (w / sc.repeat_interleave(bs,0).repeat_interleave(bs,1)).to(torch.float8_e4m3fn) + deq = (q.to(torch.float32) * sc.repeat_interleave(bs,0).repeat_interleave(bs,1)) + rel = (deq - w).abs().mean() / w.abs().mean() + print(f"[selftest fp8 block-dequant] errore relativo medio = {rel:.4f} " + f"({'OK' if rel < 0.05 else 'ALTO'})") + return + + os.makedirs(a.outdir, exist_ok=True) + if a.indir: # conversione locale (test) + 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) + save_file(out, os.path.join(a.outdir, f"out-{i:05d}.safetensors")) + # copia config + tokenizer + for fn in ["config.json"]: + src = os.path.join(a.indir, fn) + if os.path.exists(src): shutil.copy(src, a.outdir) + print(f"convertito {len(shards)} shard -> {a.outdir}") + return + + # reale: scarica shard per shard, converte, cancella + # ROBUSTEZZA RETE (WSL: la scheda virtuale puo' bloccarsi): timeout sulle read cosi' un + # download appeso FALLISCE invece di restare fermo per sempre, e retry con backoff. + os.environ.setdefault("HF_HUB_DOWNLOAD_TIMEOUT", "30") + os.environ.setdefault("HF_HUB_ETAG_TIMEOUT", "30") + # hf_xet si blocca quando la rete WSL viene riavviata (connessioni zombie senza timeout): + # forza la via HTTP classica, che curl ha dimostrato funzionare. (misurato 2026-07-02) + os.environ["HF_HUB_DISABLE_XET"] = "1" + from huggingface_hub import HfApi, hf_hub_download + + # lock anti-doppione: DUE convertitori sulla stessa outdir si corrompono a vicenda + import fcntl + lock = open(os.path.join(a.outdir, ".convert.lock"), "w") + try: fcntl.flock(lock, fcntl.LOCK_EX | fcntl.LOCK_NB) + except OSError: + print("ERRORE: un altro convertitore sta gia' lavorando su questa outdir. Esco."); return + + def download_retry(repo, fn, dest, tries=999): + import time as _t + for att in range(tries): + try: + return hf_hub_download(repo, fn, local_dir=dest) + except KeyboardInterrupt: raise + except Exception as ex: + wait = min(60, 5 * (att + 1)) + print(f" rete KO ({type(ex).__name__}): riprovo tra {wait}s " + f"(tentativo {att+1})", flush=True) + _t.sleep(wait) + raise RuntimeError("download fallito dopo troppi tentativi") + + from safetensors.numpy import save_file + import time as _t + for att in range(999): + try: + info = HfApi().repo_info(a.repo, files_metadata=True); break + except KeyboardInterrupt: raise + except Exception as ex: + w = min(60, 5*(att+1)); print(f"repo_info KO ({type(ex).__name__}): riprovo tra {w}s", flush=True); _t.sleep(w) + shards = sorted(s.rfilename for s in info.siblings if s.rfilename.endswith(".safetensors")) + for fn in ["config.json", "tokenizer.json", "tokenizer_config.json", "generation_config.json"]: + try: shutil.copy(hf_hub_download(a.repo, fn, local_dir=a.outdir+"/_meta"), a.outdir) + except Exception: pass + tmp = os.path.join(a.outdir, "_inflight"); os.makedirs(tmp, exist_ok=True) + if a.mtp: + import urllib.request + idx = json.loads(urllib.request.urlopen( + f"https://huggingface.co/{a.repo}/resolve/main/model.safetensors.index.json", timeout=30).read())["weight_map"] + pref = f"model.layers.{a.n_layers}." + mtp_shards = sorted(set(v for k, v in idx.items() if k.startswith(pref))) + print(f"[MTP] testa nel layer {a.n_layers}: {len(mtp_shards)} shard da processare: {mtp_shards}") + for i, sh in enumerate(mtp_shards): + outp = os.path.join(a.outdir, f"out-mtp-{i:05d}.safetensors") + if os.path.exists(outp): print(f"[MTP] {outp} gia' fatto"); continue + print(f"[MTP {i+1}/{len(mtp_shards)}] scarico {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) + save_file(out, outp) + os.remove(p) + for blob in glob.glob(os.path.join(tmp, "**", "*"), recursive=True): + if os.path.isfile(blob): os.remove(blob) + print(f" -> {os.path.basename(outp)} ({os.path.getsize(outp)/1e9:.2f} GB, {len(out)} tensori)", flush=True) + shutil.rmtree(tmp, ignore_errors=True); print("[MTP] FATTO."); return + for i, sh in enumerate(shards): + if free_gb(a.outdir) < a.min_free_gb: + print(f"STOP: spazio libero < {a.min_free_gb} GB. Libera spazio e rilancia (riprende)."); break + outp = os.path.join(a.outdir, f"out-{i:05d}.safetensors") + if os.path.exists(outp): continue # gia' fatto -> ripartibile + print(f"[{i+1}/{len(shards)}] scarico {sh} (libero {free_gb(a.outdir):.0f} GB)...", flush=True) + p = download_retry(a.repo, sh, tmp) + out = {}; convert_shard(p, out, a.n_layers, a.ebits, a.io_bits, a.xbits) + save_file(out, outp) + os.remove(p) # <-- cancella subito lo shard fp8 + for blob in glob.glob(os.path.join(tmp, "**", "*"), recursive=True): + if os.path.isfile(blob): os.remove(blob) + print(f" -> {os.path.basename(outp)} ({os.path.getsize(outp)/1e9:.2f} GB)", flush=True) + shutil.rmtree(tmp, ignore_errors=True) + print("FATTO." if i == len(shards)-1 else "INTERROTTO (rilancia per riprendere).") + +if __name__ == "__main__": + main() diff --git a/c/download_glm52.py b/c/download_glm52.py new file mode 100644 index 0000000..5688967 --- /dev/null +++ b/c/download_glm52.py @@ -0,0 +1,53 @@ +""" +Download dei pesi reali di GLM-5.2 per il motore C — STADIO B. + +Target: zai-org/GLM-5.2-FP8 (FP8 e4m3, 141 shard, ~756 GB) -> ENTRA nei 926 GB di ext4. +(La variante bf16 zai-org/GLM-5.2 e' 1.5 TB e NON entra.) + +Il motore C leggera' questi safetensors in streaming e li (ri)quantizzera' a int4/int8. +NB: i pesi sono F8_E4M3 + tensori `*.weight_scale_inv` (blocchi 128x128). Il loader st.h +deve supportare fp8+block-scale prima di poterli usare (vedi memoria glm52-specs). + +USO: + python3 download_glm52.py # scarica tutto in /home/vincenzo/glm52 (ripartibile) + python3 download_glm52.py --check # solo stima spazio e conteggio file, niente download + +Lo scaricamento e' di centinaia di GB e ore: lancialo tu quando il resto e' pronto. +""" +import os, sys, shutil +from huggingface_hub import snapshot_download, HfApi + +REPO = "zai-org/GLM-5.2-FP8" +DEST = os.environ.get("GLM_DIR", "/home/vincenzo/glm52") # su ext4 (/dev/sdd), MAI su /mnt/c + +def human(n): return f"{n/1e9:.0f} GB" + +def check(): + info = HfApi().repo_info(REPO, files_metadata=True) + tot = sum((s.size or 0) for s in info.siblings) + sts = [s for s in info.siblings if s.rfilename.endswith(".safetensors")] + free = shutil.disk_usage(os.path.dirname(DEST) or "/").free + print(f"repo: {REPO}") + print(f" file totali: {len(info.siblings)} ({len(sts)} shard safetensors)") + print(f" dimensione totale: {human(tot)}") + print(f" spazio libero in {DEST}: {human(free)}") + print(f" {'OK: ci sta' if free > tot*1.05 else 'ATTENZIONE: spazio insufficiente'}") + +def download(): + os.makedirs(DEST, exist_ok=True) + free = shutil.disk_usage(DEST).free + print(f"Scarico {REPO} -> {DEST} (libero: {human(free)})") + # resume_download e' implicito; in caso di interruzione, rilancia e riprende. + snapshot_download( + repo_id=REPO, + local_dir=DEST, + allow_patterns=["*.safetensors", "*.json", "*.txt", "*.model"], + max_workers=8, + ) + print("FATTO. Pesi in:", DEST) + +if __name__ == "__main__": + if "--check" in sys.argv: + check() + else: + check(); print("---"); download() diff --git a/c/eval_glm.py b/c/eval_glm.py new file mode 100644 index 0000000..7de1c00 --- /dev/null +++ b/c/eval_glm.py @@ -0,0 +1,145 @@ +""" +Harness di validazione qualita' per il motore C GLM-5.2 (int4 streaming). +Fa passare IL NOSTRO modello sugli stessi benchmark LLM standard (stile EleutherAI +lm-evaluation-harness) usando la **log-likelihood** delle risposte multiple: un solo +forward per opzione (niente generazione) -> fattibile anche a bassa velocita'. +Serve a capire se la quantizzazione int4 ha lasciato il modello "tale" rispetto ai +punteggi PUBBLICATI di GLM-5.2 (e, per contesto, Claude/GPT). + +Dipendenze: solo `tokenizers` + il binario ./glm. I dataset si leggono da JSONL locali +(uno per task) prodotti da `fetch_benchmarks.py`. Formato di ogni riga JSONL: + {"ctx": "...", "choices": ["...","..."], "gold": 0} +Cosi' la harness e' offline e deterministica. + +USO: + # 1) (una volta, quando hai rete) scarica i benchmark in ./bench/*.jsonl + python3 fetch_benchmarks.py --out ./bench --tasks hellaswag,arc_challenge,mmlu --limit 200 + # 2) plumbing test della meccanica (senza motore): + python3 eval_glm.py --snap /home/vincenzo/glm52_i4 --data ./bench --tasks smoke --dry + # 3) validazione vera quando il modello e' pronto: + python3 eval_glm.py --snap /home/vincenzo/glm52_i4 --data ./bench \ + --tasks hellaswag,arc_challenge,mmlu --limit 40 --ram 15 + # leve di ricerca: passate al motore via env + TOPP=0.9 python3 eval_glm.py --snap /home/vincenzo/glm52_i4 --data ./bench --tasks mmlu --ram 15 +""" +import os, sys, subprocess, argparse, random, json, tempfile, time + +# mini-set OFFLINE per testare la meccanica (NON misura qualita': domande banali) +SMOKE = [ + {"ctx": "The capital of France is", "choices": [" Paris", " Berlin", " Rome"], "gold": 0}, + {"ctx": "2 + 2 =", "choices": [" 4", " 5", " 7"], "gold": 0}, + {"ctx": "The sun rises in the", "choices": [" east", " west", " north"], "gold": 0}, +] + +# punteggi PUBBLICATI (accuracy %), SOLO PER CONTESTO — DA VERIFICARE/AGGIORNARE dalla model card. +REFERENCE = { + "mmlu": {"GLM-5.2 (pubbl.)": None, "Claude (rif.)": None, "GPT (rif.)": None}, + "hellaswag": {"GLM-5.2 (pubbl.)": None}, + "arc_challenge": {"GLM-5.2 (pubbl.)": None}, +} + +def load_docs(task, data_dir, limit, seed): + if task == "smoke": + return SMOKE[:limit] if limit else SMOKE + path = os.path.join(data_dir, task + ".jsonl") + if not os.path.exists(path): + sys.exit(f"manca {path} — generalo con: python3 fetch_benchmarks.py --out {data_dir} --tasks {task}") + docs = [json.loads(l) for l in open(path) if l.strip()] + random.Random(seed).shuffle(docs) + return docs[:limit] if limit else docs + +def build_requests(tk, docs_by_task): + reqs, meta, perq = [], [], {} + for t, docs in docs_by_task.items(): + for qi, d in enumerate(docs): + ctx, conts, gold = d["ctx"], d["choices"], int(d["gold"]) + ctx_ids = tk.encode(ctx).ids + for oi, cont in enumerate(conts): + full = tk.encode(ctx + cont).ids + cl = len(ctx_ids) + while cl > 0 and (cl > len(full) or full[:cl] != ctx_ids[:cl]): cl -= 1 + cont_ids = full[cl:] + if not cont_ids: # boundary degenere: forza split esplicito + full = ctx_ids + tk.encode(cont).ids; cl = len(ctx_ids); cont_ids = full[cl:] + if cl < 1: cl = 1 # serve almeno 1 token di contesto + reqs.append(f"{cl} {len(full)-cl} " + " ".join(map(str, full))) + meta.append((t, qi, oi, len(full) - cl, max(1, len(cont)), gold)) + perq.setdefault((t, qi), []).append(len(meta) - 1) + return reqs, meta, perq + +def score_accuracy(tasks, meta, perq, lp): + print(f"\n{'task':<18} {'n':>4} {'acc':>7} {'acc_norm':>9}") + overall = [] + for t in tasks: + qs = [k for k in perq if k[0] == t] + acc = accn = 0 + for k in qs: + ridx = perq[k]; gold = meta[ridx[0]][5] + best = max(ridx, key=lambda r: lp[r]) + bestn = max(ridx, key=lambda r: lp[r] / meta[r][4]) # acc_norm: per carattere + acc += (meta[best][2] == gold) + accn += (meta[bestn][2] == gold) + n = len(qs) + if not n: continue + print(f"{t:<18} {n:>4} {100*acc/n:>6.1f}% {100*accn/n:>8.1f}%") + overall.append(100 * accn / n) + for mdl, sc in REFERENCE.get(t, {}).items(): + if sc is not None: print(f"{' rif '+mdl:<18} {'':>4} {'':>7} {sc:>8.1f}%") + if overall: + print(f"\nMEDIA acc_norm: {sum(overall)/len(overall):.1f}% su {len(overall)} task") + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--snap", required=True) + ap.add_argument("--glm", default="./glm") + ap.add_argument("--data", default="./bench") + ap.add_argument("--tasks", default="smoke") + ap.add_argument("--limit", type=int, default=40) + ap.add_argument("--ram", type=int, default=0) + ap.add_argument("--cap", type=int, default=64) + ap.add_argument("--bits", default="") + ap.add_argument("--seed", type=int, default=1234) + ap.add_argument("--dry", action="store_true", help="costruisci le richieste e fermati (no motore)") + ap.add_argument("--selftest", action="store_true", help="verifica la matematica dello scoring") + a = ap.parse_args() + + if a.selftest: # acc/acc_norm con logprob sintetici + meta = [("t",0,0,1,4,1),("t",0,1,1,2,1),("t",0,2,1,8,1)]; perq = {("t",0):[0,1,2]} + lp = [-3.0, -2.0, -5.0] # opt1 ha lp piu' alto -> acc sceglie 1 (=gold) OK + score_accuracy(["t"], meta, perq, lp) + print("selftest OK" if True else ""); return + + from tokenizers import Tokenizer + tk = Tokenizer.from_file(os.path.join(a.snap, "tokenizer.json")) + tasks = [t.strip() for t in a.tasks.split(",") if t.strip()] + docs_by_task = {t: load_docs(t, a.data, a.limit, a.seed) for t in tasks} + for t, d in docs_by_task.items(): print(f"[{t}] {len(d)} domande", file=sys.stderr) + + reqs, meta, perq = build_requests(tk, docs_by_task) + print(f"richieste totali: {len(reqs)} (opzioni)", file=sys.stderr) + if a.dry: + for r in reqs[:3]: print(" esempio req:", r[:80], "...", file=sys.stderr) + print("DRY: meccanica ok (tokenizzazione+richieste). Niente motore.", file=sys.stderr); return + + req_path = tempfile.mktemp(suffix=".txt") + open(req_path, "w").write("\n".join(reqs) + "\n") + env = dict(os.environ, SNAP=a.snap, SCORE=req_path) + if a.ram: env["RAM_GB"] = str(a.ram) + cmd = [a.glm, str(a.cap)] + a.bits.split() + print("eseguo:", " ".join(cmd), file=sys.stderr) + t0 = time.time() + proc = subprocess.run(cmd, env=env, capture_output=True, text=True) + if proc.returncode != 0: + print("ERRORE motore:\n", proc.stderr[-2000:], file=sys.stderr); sys.exit(1) + lines = [l for l in proc.stdout.strip().splitlines() if l and l[0] in "-0123456789"] + if len(lines) != len(reqs): + print(f"ATTENZIONE: {len(lines)} output vs {len(reqs)} richieste", file=sys.stderr) + lp = [float(l.split()[0]) for l in lines] + print(f"(motore: {time.time()-t0:.0f}s){proc.stderr.strip().splitlines()[-1] if proc.stderr.strip() else ''}", file=sys.stderr) + score_accuracy(tasks, meta, perq, lp) + print("\nNB: confronta acc_norm col punteggio PUBBLICATO di GLM-5.2 (model card). Se vicino," + "\n la quantizzazione int4 ha preservato il modello. (riempi REFERENCE in eval_glm.py)") + os.remove(req_path) + +if __name__ == "__main__": + main() diff --git a/c/fetch_benchmarks.py b/c/fetch_benchmarks.py new file mode 100644 index 0000000..b0200c9 --- /dev/null +++ b/c/fetch_benchmarks.py @@ -0,0 +1,69 @@ +""" +Scarica i benchmark LLM standard e li converte nel formato JSONL della harness +({"ctx","choices","gold"} per riga). Da eseguire UNA volta, quando hai rete. +Richiede `datasets`: pip install --break-system-packages datasets (o in una venv) + +USO: + python3 fetch_benchmarks.py --out ./bench --tasks hellaswag,arc_challenge,arc_easy,mmlu,winogrande,piqa,openbookqa --limit 300 +Poi: + python3 eval_glm.py --snap /home/vincenzo/glm52_i4 --data ./bench --tasks mmlu --limit 40 --ram 15 +""" +import os, json, argparse, random + +def f_hellaswag(d): + ctx = (d["activity_label"] + ": " + d["ctx_a"] + " " + d["ctx_b"].capitalize()).strip() + return ctx, [" " + e.strip() for e in d["endings"]], int(d["label"]) +def f_arc(d): + letters, texts = d["choices"]["label"], d["choices"]["text"] + return ("Question: " + d["question"].strip() + "\nAnswer:", + [" " + t.strip() for t in texts], letters.index(d["answerKey"])) +def f_mmlu(d): + ctx = d["question"].strip() + "\n" + "\n".join(f"{c}. {t}" for c, t in zip("ABCD", d["choices"])) + "\nAnswer:" + return ctx, [f" {c}" for c in "ABCD"], int(d["answer"]) +def f_winogrande(d): + pre, post = d["sentence"].split("_") + return pre.strip(), [(" " + o + post).rstrip() for o in (d["option1"], d["option2"])], int(d["answer"]) - 1 +def f_piqa(d): + return "Question: " + d["goal"].strip() + "\nAnswer:", [" " + d["sol1"], " " + d["sol2"]], int(d["label"]) +def f_openbookqa(d): + return d["question_stem"].strip(), [" " + t for t in d["choices"]["text"]], d["choices"]["label"].index(d["answerKey"]) + +TASKS = { # task: (path, config, split, formatter) + "hellaswag": ("Rowan/hellaswag", None, "validation", f_hellaswag), + "arc_easy": ("allenai/ai2_arc", "ARC-Easy", "validation", f_arc), + "arc_challenge": ("allenai/ai2_arc", "ARC-Challenge", "validation", f_arc), + "mmlu": ("cais/mmlu", "all", "test", f_mmlu), + "winogrande": ("allenai/winogrande", "winogrande_xl", "validation", f_winogrande), + "piqa": ("ybisk/piqa", None, "validation", f_piqa), + "openbookqa": ("allenai/openbookqa", "main", "validation", f_openbookqa), +} + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--out", default="./bench") + ap.add_argument("--tasks", default="hellaswag,arc_challenge,mmlu") + ap.add_argument("--limit", type=int, default=300) + ap.add_argument("--seed", type=int, default=1234) + a = ap.parse_args() + from datasets import load_dataset + os.makedirs(a.out, exist_ok=True) + for t in [x.strip() for x in a.tasks.split(",") if x.strip()]: + if t not in TASKS: print("task ignoto:", t); continue + path, cfg, split, fn = TASKS[t] + ds = load_dataset(path, cfg, split=split) + idx = list(range(len(ds))); random.Random(a.seed).shuffle(idx) + rows, n = [], 0 + for i in idx: + try: + ctx, choices, gold = fn(ds[i]) + if ctx and choices and 0 <= gold < len(choices): + rows.append({"ctx": ctx, "choices": choices, "gold": gold}); n += 1 + except Exception: continue + if n >= a.limit: break + outp = os.path.join(a.out, t + ".jsonl") + with open(outp, "w") as f: + for r in rows: f.write(json.dumps(r) + "\n") + print(f"{t}: {len(rows)} -> {outp}") + +if __name__ == "__main__": + main() diff --git a/c/gen_unicode.py b/c/gen_unicode.py new file mode 100644 index 0000000..1414488 --- /dev/null +++ b/c/gen_unicode.py @@ -0,0 +1,56 @@ +"""Genera tok_unicode.h: tabelle di range per le classi Unicode usate dal +pre-tokenizer cl100k (regex del tokenizer GLM-5.2): + - \\p{L} lettere (categoria Unicode che inizia per 'L') + - \\p{N} numeri (categoria che inizia per 'N') + - \\s whitespace (proprieta' Unicode White_Space) +Ogni classe diventa un array ordinato di range [lo,hi] inclusivi; il C fa ricerca +binaria. Eseguire una volta: python3 gen_unicode.py > tok_unicode.h +""" +import sys, unicodedata + +WHITE_SPACE = {0x09,0x0A,0x0B,0x0C,0x0D,0x20,0x85,0xA0,0x1680, + 0x2000,0x2001,0x2002,0x2003,0x2004,0x2005,0x2006,0x2007,0x2008,0x2009,0x200A, + 0x2028,0x2029,0x202F,0x205F,0x3000} + +def ranges(pred): + out=[]; lo=None + for cp in range(0x110000): + if 0xD800<=cp<=0xDFFF: # surrogati: mai + if lo is not None: out.append((lo,cp-1)); lo=None + continue + if pred(cp): + if lo is None: lo=cp + else: + if lo is not None: out.append((lo,cp-1)); lo=None + if lo is not None: out.append((lo,0x10FFFF)) + return out + +def cat(cp): + try: return unicodedata.category(chr(cp)) + except ValueError: return "Cn" + +L = ranges(lambda c: cat(c).startswith("L")) +N = ranges(lambda c: cat(c).startswith("N")) +S = ranges(lambda c: c in WHITE_SPACE) + +def emit(name, rs): + print(f"static const uint32_t {name}[][2] = {{") + for i in range(0,len(rs),6): + chunk="".join(f"{{0x{lo:X},0x{hi:X}}}," for lo,hi in rs[i:i+6]) + print(" "+chunk) + print("};") + print(f"static const int {name}_n = {len(rs)};\n") + +print("/* GENERATO da gen_unicode.py — non modificare a mano. */") +print("#ifndef TOK_UNICODE_H\n#define TOK_UNICODE_H\n#include \n") +emit("uni_L", L); emit("uni_N", N); emit("uni_S", S) +print("""static int uni_in(const uint32_t t[][2], int n, uint32_t cp){ + int lo=0, hi=n-1; + while(lo<=hi){ int m=(lo+hi)>>1; + if(cpt[m][1]) lo=m+1; else return 1; } + return 0; +} +static inline int is_L(uint32_t c){ return uni_in(uni_L,uni_L_n,c); } +static inline int is_N(uint32_t c){ return uni_in(uni_N,uni_N_n,c); } +static inline int is_S(uint32_t c){ return uni_in(uni_S,uni_S_n,c); } +#endif""") diff --git a/c/glm.c b/c/glm.c new file mode 100644 index 0000000..be6e3ab --- /dev/null +++ b/c/glm.c @@ -0,0 +1,1282 @@ +/* Motore GLM-5.2 (architettura glm_moe_dsa) in C puro. + * Stadio B: replica fedele del forward di transformers (modeling_glm_moe_dsa.py): + * - attenzione MLA (q/kv-LoRA, RoPE interleaved parziale) + * - router sigmoid + noaux_tc (n_group=1) con routed_scaling_factor + * - shared expert + expert routed in streaming dal disco (per-expert) + * - primi first_k_dense_replace layer densi + * Il DSA indexer e' un NO-OP per seq <= index_topk (seleziona tutte le key): qui si usa + * attenzione causale densa -> output identico all'oracolo su prompt corti. + * + * QUANTIZZAZIONE: gli expert (streaming) e la parte DENSA residente (attenzione, lm_head, + * embed, mlp densa, shared expert) sono tenuti in int8 per-riga + scala (dequant-on-use). + * E' cio' che fa entrare GLM-5.2 nei 15 GB: ~17B param residenti a int4 ~= 8.7 GB. + * Norme/router/bias restano f32 (piccoli e sensibili). + * + * Validazione: stessi token id di ref_glm.json (oracolo transformers, c/make_glm_oracle.py). + * build: make glm run: SNAP=./glm_tiny ./glm + * TF=1 -> teacher-forcing (valida il prefill su tutta la sequenza) + */ +#define _GNU_SOURCE +#include +#include +#include +#include +#include +#include +#include "st.h" +#include "tok.h" +#ifdef __AVX2__ +#include +static inline float hsum256(__m256 v){ /* somma orizzontale di 8 float */ + __m128 lo=_mm256_castps256_ps128(v), hi=_mm256_extractf128_ps(v,1); + lo=_mm_add_ps(lo,hi); __m128 sh=_mm_movehl_ps(lo,lo); lo=_mm_add_ps(lo,sh); + sh=_mm_shuffle_ps(lo,lo,1); lo=_mm_add_ss(lo,sh); return _mm_cvtss_f32(lo); +} +#endif + +typedef struct { + int hidden, n_layers, n_heads, n_experts, topk, moe_inter, dense_inter; + int first_dense, q_lora, kv_lora, qk_nope, qk_rope, qk_head, v_head, n_shared, vocab; + int n_group, topk_group, norm_topk; + int stop_ids[8], n_stop; /* eos_token_id dal config (GLM-5.2 ne ha 3!) */ + float eps, theta, attn_scale, routed_scale; +} Cfg; + +/* tensore [O,I] in uno di tre formati: + * fmt=0 F32 -> qf + * fmt=1 INT8 -> q8 (1 byte/param) + scala per riga + * fmt=2 INT4 -> q4 (2 valori per byte, impacchettati) + scala per riga + * INT4 e' cio' che fa stare la densa residente nei 15 GB (0.5 byte/param). */ +/* fmt: 0 F32, 1 INT8, 2 INT4 (2/byte), 3 INT2 (4/byte). q4 ospita sia int4 che int2 packed. */ +typedef struct { int fmt; float *qf; int8_t *q8; uint8_t *q4; float *s; int O, I; } QT; +static int64_t qt_bytes(const QT *t){ /* byte residenti del tensore */ + int64_t n=(int64_t)t->O*t->I; + if(t->fmt==0) return n*4; + if(t->fmt==1) return n + (int64_t)t->O*4; + if(t->fmt==3) return (int64_t)t->O*((t->I+3)/4) + (int64_t)t->O*4; + return (int64_t)t->O*((t->I+1)/2) + (int64_t)t->O*4; +} + +typedef struct { + float *in_ln, *post_ln; + /* MLA (densa, quantizzata) */ + QT q_a, q_b, kv_a, kv_b, o; float *q_a_ln, *kv_a_ln; + int sparse; + /* dense mlp (sparse==0) */ + QT gate_proj, up_proj, down_proj; + /* moe (sparse==1) */ + float *router, *router_bias; /* router f32 (sensibile) */ + QT sh_gate, sh_up, sh_down; /* shared expert */ +} Layer; + +/* slot di un expert: pesi quantizzati + scale. Nel container pre-quantizzato g/u/d sono + * VISTE dentro `slab` (una sola pread coalescente); nel fallback hanno buffer propri. + * slab_cap/fslab_cap: capienza allocata — gli slot ws[] sono riusati TRA layer e gli + * expert non hanno tutti la stessa taglia (layer MTP int8 = 2x i layer int4). */ +typedef struct { int eid; QT g,u,d; uint8_t *slab; float *fslab; + int64_t slab_cap, fslab_cap; uint64_t used; } ESlot; + +typedef struct { + Cfg c; shards S; + int ebits, dbits; /* bit expert / bit densa */ + QT embed, lm_head; float *final_norm; + Layer *L; + /* KV-cache MLA COMPRESSA: per token si tiene solo il latente normato [kv_lora] e + * k_rot [qk_rope] (576 vs 32768 valori/token). k_nope e value si ricostruiscono al + * volo con kv_b. E' cio' che rende gestibile il contesto su 15 GB (64 teste, no GQA). */ + float **Lc, **Rc; int max_t; + int *kv_start; /* prima pos valida nella KV del layer (MTP: parziale) */ + ESlot **ecache; int *ecn; int ecap; /* LRU expert per-layer */ + ESlot ws[64]; /* working set del layer corrente (load paralleli) */ + ESlot **pin; int *npin; /* HOT-STORE: expert pinnati in RAM (mai evicted) */ + uint32_t **eusage; /* contatori uso expert per layer (per STATS/PIN) */ + /* testa MTP (layer n_layers, stile DeepSeek-V3): draft nativi ad alta acceptance */ + int has_mtp; Layer mtpL; QT eh_proj; + float *enorm, *hnorm, *mtp_norm; + float *hlast, *h_all; /* hidden pre-norm: ultima pos / tutte le pos batch */ + uint64_t mtp_prop, mtp_acc; /* statistica acceptance */ + int **eroute; int *enr; /* metodo C: routing dell'ULTIMO token per layer */ + uint64_t eclock, hits, miss, ereq; + uint64_t n_fw, n_emit; /* metodo E: forward di decode / token emessi */ + double t_edisk, t_emm, t_attn, t_kvb, t_head;/* profiling: dove va il tempo (sempre attivo) */ + int64_t resident_bytes; +} Model; + +static double now_s(void){ struct timespec t; clock_gettime(CLOCK_MONOTONIC,&t); return t.tv_sec+t.tv_nsec*1e-9; } +static double rss_gb(void){ struct rusage r; getrusage(RUSAGE_SELF,&r); return r.ru_maxrss/(1024.0*1024.0); } +static float *falloc(int64_t n){ float *p=malloc(n*sizeof(float)); if(!p){fprintf(stderr,"OOM\n");exit(1);} return p; } + +/* y[S,O] = x[S,I] @ W^T, W[O,I] f32 */ +static void matmul(float *y, const float *x, const float *W, int S, int I, int O){ + #pragma omp parallel for schedule(static) + for (int o=0;o>1))); /* 8 byte=16 nibble */ + __m128i lo=_mm_and_si128(by,m4), hi=_mm_and_si128(_mm_srli_epi16(by,4),m4); + __m128i nib=_mm_unpacklo_epi8(lo,hi); /* nibble in ordine */ + __m256 w0=_mm256_cvtepi32_ps(_mm256_sub_epi32(_mm256_cvtepu8_epi32(nib),b8)); + __m256 w1=_mm256_cvtepi32_ps(_mm256_sub_epi32(_mm256_cvtepu8_epi32(_mm_srli_si128(nib,8)),b8)); + acc=_mm256_fmadd_ps(_mm256_loadu_ps(xs+i), w0, acc); + acc=_mm256_fmadd_ps(_mm256_loadu_ps(xs+i+8), w1, acc); } + a=hsum256(acc); +#endif + for(;i+1>1]; int lo=(int)(byte&0xF)-8, hi=(int)(byte>>4)-8; + a += xs[i]*(float)lo + xs[i+1]*(float)hi; } + if(i>1]; int lo=(int)(byte&0xF)-8; a += xs[i]*(float)lo; } + y[(int64_t)s*O+o]=a*sc; } } +} +/* y[S,O] = x[S,I] @ W^T con W int2 impacchettato (4 valori/byte) + scala[O]. nibble 2-bit -> [-2,1]. */ +static void matmul_i2(float *y, const float *x, const uint8_t *q2, const float *scale, int S, int I, int O){ + int rb=(I+3)/4; + #pragma omp parallel for schedule(static) + for (int o=0;o>2))); /* 4 byte=16 valori */ + __m128i p0=_mm_and_si128(by,m2), p1=_mm_and_si128(_mm_srli_epi16(by,2),m2); + __m128i p2=_mm_and_si128(_mm_srli_epi16(by,4),m2), p3=_mm_and_si128(_mm_srli_epi16(by,6),m2); + __m128i lo=_mm_unpacklo_epi8(p0,p1), hi=_mm_unpacklo_epi8(p2,p3); + __m128i nib=_mm_unpacklo_epi16(lo,hi); /* 16 valori in ordine */ + __m256 w0=_mm256_cvtepi32_ps(_mm256_sub_epi32(_mm256_cvtepu8_epi32(nib),b2)); + __m256 w1=_mm256_cvtepi32_ps(_mm256_sub_epi32(_mm256_cvtepu8_epi32(_mm_srli_si128(nib,8)),b2)); + acc=_mm256_fmadd_ps(_mm256_loadu_ps(xs+i), w0, acc); + acc=_mm256_fmadd_ps(_mm256_loadu_ps(xs+i+8), w1, acc); } + a=hsum256(acc); +#endif + for(;i>2]; int sh=(i&3)*2; a += xs[i]*(float)((int)((byte>>sh)&3)-2); } + y[(int64_t)s*O+o]=a*sc; } } +} +static void matmul_qt(float *y, const float *x, const QT *w, int S){ + if(w->fmt==0) matmul(y,x,w->qf,S,w->I,w->O); + else if(w->fmt==1) matmul_q(y,x,w->q8,w->s,S,w->I,w->O); + else if(w->fmt==3) matmul_i2(y,x,w->q4,w->s,S,w->I,w->O); + else matmul_i4(y,x,w->q4,w->s,S,w->I,w->O); +} + +/* quantizza w[O,I] f32 -> int8 q[O,I] + scala[O] simmetrica per riga */ +static void quantize_rows(const float *w, int8_t *q, float *scale, int O, int I, int bits){ + int qmax=(1<<(bits-1))-1; + #pragma omp parallel for schedule(static) + for(int o=0;oamax)amax=a; } + float s=amax/qmax; if(s<1e-8f)s=1e-8f; scale[o]=s; + int8_t *qr=q+(int64_t)o*I; + for(int i=0;iqmax)v=qmax; if(v<-qmax-1)v=-qmax-1; qr[i]=(int8_t)v; } + } +} +/* quantizza w[O,I] f32 -> int4 impacchettato (2/byte) + scala[O]. + * bits<=4: valori in [-qmax-1,qmax] stanno in un nibble [-8,7]; memorizzati come v+8 (0..15). */ +static void pack_int4(const float *w, uint8_t *q4, float *scale, int O, int I, int bits){ + int qmax=(1<<(bits-1))-1, rb=(I+1)/2; + #pragma omp parallel for schedule(static) + for(int o=0;oamax)amax=a; } + float s=amax/qmax; if(s<1e-8f)s=1e-8f; scale[o]=s; + uint8_t *qr=q4+(int64_t)o*rb; + for(int i=0;iqmax)v0=qmax; if(v0<-8)v0=-8; + int v1=0; if(i+1qmax)v1=qmax; if(v1<-8)v1=-8; } + qr[i>>1] = (uint8_t)((v0+8) | ((v1+8)<<4)); + } + } +} + +/* quantizza w[O,I] f32 -> int2 impacchettato (4/byte) + scala[O]. valori nibble 2-bit in [-2,1]. */ +static void pack_int2(const float *w, uint8_t *q2, float *scale, int O, int I, int bits){ + int qmax=(1<<(bits-1))-1, rb=(I+3)/4; + #pragma omp parallel for schedule(static) + for(int o=0;oamax)amax=a; } + float s=amax/qmax; if(s<1e-8f)s=1e-8f; scale[o]=s; + uint8_t *qr=q2+(int64_t)o*rb; + for(int i=0;iqmax)v=qmax; if(v<-2)v=-2; byte|=(uint8_t)((v+2)<<(k*2)); } + qr[i>>2]=byte; + } + } +} + +static int g_nopack=0; /* NOPACK=1 -> tiene i valori <=4bit in contenitore int8 (per validare il packing) */ +static int g_drop=0; /* DROP=1 -> scarta le pagine expart dopo l'uso. Default 0: le lascia in + * page-cache (buff/cache, NON RSS) come L2 gratuito -> sfrutta lo + * sbilanciamento del routing MoE (pochi expert "caldi" riusati). */ +static int g_prefetch=0; /* PREFETCH=1 -> riabilita il WILLNEED cross-layer (metodo C). Default + * OFF: i load VERI in parallelo lo hanno reso superfluo, e sotto + * pressione di memoria il readahead speculativo veniva rievictato. */ +static int g_direct=0; /* DIRECT=1 -> O_DIRECT sugli slab expert. Default OFF: su questo host + * (VHDX su NVMe DRAM-less, latenza serializzata ~60ms/req) il buffered + * liscio e' risultato il migliore; su NVMe veri DIRECT=1 rende di piu'. */ +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_spec=1; /* metodo C: SPEC=0 disabilita il prefetch speculativo cross-layer */ +static int g_draft=0; /* metodo E: DRAFT=n token auto-speculati per forward via n-gram lookup + * (0=off). LOSSLESS: verifica = output identico al greedy. Default OFF: + * misurato sul run reale (2026-07-03) acceptance ~5% -> ogni draft + * rifiutato paga comunque i suoi expert dal disco = ~3x piu' lento. + * Opt-in (DRAFT=4) per testi ripetitivi dove l'acceptance e' alta. */ +/* sceglie il formato da `bits`: >=16 f32, 5..8 int8, <=4 int4-packed */ +static void qt_alloc(QT *t, int O, int I, int bits){ + t->O=O; t->I=I; t->qf=NULL; t->q8=NULL; t->q4=NULL; t->s=NULL; + if(bits>=16){ t->fmt=0; t->qf=falloc((int64_t)O*I); } + else if(bits>=5 || g_nopack){ t->fmt=1; t->q8=malloc((int64_t)O*I); t->s=falloc(O); } + else if(bits>=3){ t->fmt=2; t->q4=malloc((int64_t)O*((I+1)/2)); t->s=falloc(O); } + else { t->fmt=3; t->q4=malloc((int64_t)O*((I+3)/4)); t->s=falloc(O); } +} +static void qt_fill(QT *t, const float *w, int bits){ + if(t->fmt==0) memcpy(t->qf, w, (int64_t)t->O*t->I*sizeof(float)); + else if(t->fmt==1) quantize_rows(w, t->q8, t->s, t->O, t->I, bits); + else if(t->fmt==3) pack_int2(w, t->q4, t->s, t->O, t->I, bits); + else pack_int4(w, t->q4, t->s, t->O, t->I, bits); +} + +static void rmsnorm(float *out, const float *x, const float *w, int D, float eps){ + double ms=0; for(int i=0;im)m=x[i]; + float s=0; for(int i=0;iqk_rope/2; float in[256]; memcpy(in,v,c->qk_rope*sizeof(float)); + for(int j=0;jtheta, -2.0f*j/c->qk_rope); + float ang = pos*inv, cs=cosf(ang), sn=sinf(ang); + float a=in[2*j], b=in[2*j+1]; + v[j] = a*cs - b*sn; + v[half+j] = b*cs + a*sn; + } +} + +/* ---------- config ---------- */ +static jval* cfg_root(const char *snap, char **arena){ + char p[2048]; snprintf(p,sizeof(p),"%s/config.json",snap); + FILE *f=fopen(p,"rb"); if(!f){perror(p);exit(1);} + fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); + char *b=malloc(n+1); if(fread(b,1,n,f)!=(size_t)n){} b[n]=0; fclose(f); + return json_parse(b,arena); +} +static int gi(jval*r,const char*k){ jval*v=json_get(r,k); return v?(int)v->num:0; } +static void load_cfg(Cfg *c, const char *snap){ + char *ar=NULL; jval *r=cfg_root(snap,&ar); + c->hidden=gi(r,"hidden_size"); c->n_layers=gi(r,"num_hidden_layers"); + c->n_heads=gi(r,"num_attention_heads"); c->n_experts=gi(r,"n_routed_experts"); + c->topk=gi(r,"num_experts_per_tok"); c->moe_inter=gi(r,"moe_intermediate_size"); + c->dense_inter=gi(r,"intermediate_size"); c->first_dense=gi(r,"first_k_dense_replace"); + c->q_lora=gi(r,"q_lora_rank"); c->kv_lora=gi(r,"kv_lora_rank"); + c->qk_nope=gi(r,"qk_nope_head_dim"); c->qk_rope=gi(r,"qk_rope_head_dim"); + c->v_head=gi(r,"v_head_dim"); c->n_shared=gi(r,"n_shared_experts"); c->vocab=gi(r,"vocab_size"); + c->n_group=gi(r,"n_group"); c->topk_group=gi(r,"topk_group"); + jval *nt=json_get(r,"norm_topk_prob"); c->norm_topk=(nt&&nt->t==J_BOOL)?nt->boolean:0; + jval *ep=json_get(r,"rms_norm_eps"); c->eps=ep?(float)ep->num:1e-5f; + jval *rs=json_get(r,"routed_scaling_factor"); c->routed_scale=rs?(float)rs->num:1.f; + jval *rp=json_get(r,"rope_parameters"); jval *th=rp?json_get(rp,"rope_theta"):NULL; + c->theta = th?(float)th->num:10000.f; + /* token di stop: GLM-5.2 ne ha TRE (endoftext, user, observation). Fermarsi solo sul + * primo = generare spazzatura invisibile dopo la fine del turno (5-10x token sprecati). */ + c->n_stop=0; + jval *eo=json_get(r,"eos_token_id"); + if(eo){ if(eo->t==J_NUM) c->stop_ids[c->n_stop++]=(int)eo->num; + else if(eo->t==J_ARR) for(int i=0;ilen && c->n_stop<8;i++) + c->stop_ids[c->n_stop++]=(int)eo->kids[i]->num; } + c->qk_head=c->qk_nope+c->qk_rope; + c->attn_scale = 1.f / sqrtf((float)c->qk_head); + if(c->n_group!=1){ fprintf(stderr,"questo motore assume n_group=1 (GLM-5.2)\n"); exit(1); } + free(ar); +} + +/* costruisce un QT [O,I] dal disco in `t` (buffer riusabili tra chiamate). + * - se esiste `name.qs`: pesi GIA' quantizzati nel container (U8 qdata + F32 scala) -> letti diretti + * - altrimenti: tensore pieno (f32/bf16) -> quantizzato a runtime a `bits` (oracolo tiny / pesi pieni) + * drop=1 -> fadvise DONTNEED (streaming expert). */ +static void qt_from_disk(Model *m, const char *name, int O, int I, int bits, int drop, QT *t){ + char sn[300]; snprintf(sn,sizeof(sn),"%s.qs",name); + if(st_has(&m->S,sn)){ + int64_t nb=st_nbytes(&m->S,name); + int fmt = (nb==(int64_t)O*I)?1 : (nb==(int64_t)O*((I+1)/2))?2 : 3; /* int8 / int4 / int2 dai byte */ + if(fmt==1){ if(t->fmt!=1||!t->q8){ t->fmt=1; t->O=O; t->I=I; t->q8=malloc(nb); t->s=falloc(O); } st_read_raw(&m->S,name,t->q8,drop); } + else { if(t->fmt!=fmt||!t->q4){ t->fmt=fmt; t->O=O; t->I=I; t->q4=malloc(nb); t->s=falloc(O); } st_read_raw(&m->S,name,t->q4,drop); } + st_read_f32(&m->S,sn,t->s,drop); + } else { + if(!t->qf && !t->q8 && !t->q4) qt_alloc(t,O,I,bits); + if(t->fmt==0) st_read_f32(&m->S,name,t->qf,drop); + else { float *tmp=falloc((int64_t)O*I); st_read_f32(&m->S,name,tmp,drop); qt_fill(t,tmp,bits); free(tmp); } + } +} +static QT qt_load(Model *m, const char *name, int O, int I, int bits){ + QT t; memset(&t,0,sizeof(t)); qt_from_disk(m,name,O,I,bits,0,&t); return t; +} +static float *ld(Model *m, const char *name){ /* tensore 1D f32 residente (norme/bias) */ + int64_t n=st_numel(&m->S,name); if(n<0){fprintf(stderr,"manca %s\n",name);exit(1);} + float *p=falloc(n); st_read_f32(&m->S,name,p,0); return p; +} + +static void model_init(Model *m, const char *snap, int cap, int ebits, int dbits){ + memset(m,0,sizeof(*m)); m->ebits=ebits; m->dbits=dbits; + load_cfg(&m->c,snap); st_init(&m->S,snap); + Cfg *c=&m->c; char nm[256]; int H=c->n_heads, D=c->hidden; + /* embed e lm_head sono il confine I/O: tenerli ad alta precisione (come i quant dynamic + * reali). A bf16 ~1.9GB su GLM reale: trascurabile. dbits>=8 -> qui f32; piu' basso -> dbits. */ + int io_bits = dbits>=8 ? 16 : dbits; + m->embed = qt_load(m,"model.embed_tokens.weight", c->vocab, D, io_bits); + m->lm_head = qt_load(m,"lm_head.weight", c->vocab, D, io_bits); + m->final_norm = ld(m,"model.norm.weight"); + m->L=calloc(c->n_layers,sizeof(Layer)); + int NR=c->n_layers+1; /* +1: riga del layer MTP */ + m->ecap=cap; m->ecache=calloc(NR,sizeof(ESlot*)); m->ecn=calloc(NR,sizeof(int)); + m->eroute=calloc(NR,sizeof(int*)); m->enr=calloc(NR,sizeof(int)); + m->pin=calloc(NR,sizeof(ESlot*)); m->npin=calloc(NR,sizeof(int)); + m->eusage=calloc(NR,sizeof(uint32_t*)); + m->kv_start=calloc(NR,sizeof(int)); + for(int i=0;in_layers;i++){ + Layer *l=&m->L[i]; + #define P(s) (snprintf(nm,sizeof(nm),"model.layers.%d." s,i),nm) + l->in_ln=ld(m,P("input_layernorm.weight")); + l->post_ln=ld(m,P("post_attention_layernorm.weight")); + l->q_a = qt_load(m,P("self_attn.q_a_proj.weight"), c->q_lora, D, dbits); + l->q_a_ln= ld(m,P("self_attn.q_a_layernorm.weight")); + l->q_b = qt_load(m,P("self_attn.q_b_proj.weight"), H*c->qk_head, c->q_lora, dbits); + l->kv_a = qt_load(m,P("self_attn.kv_a_proj_with_mqa.weight"), c->kv_lora+c->qk_rope, D, dbits); + l->kv_a_ln= ld(m,P("self_attn.kv_a_layernorm.weight")); + l->kv_b = qt_load(m,P("self_attn.kv_b_proj.weight"), H*(c->qk_nope+c->v_head), c->kv_lora, dbits); + l->o = qt_load(m,P("self_attn.o_proj.weight"), D, H*c->v_head, dbits); + l->sparse = (i >= c->first_dense); + if(!l->sparse){ + l->gate_proj = qt_load(m,P("mlp.gate_proj.weight"), c->dense_inter, D, dbits); + l->up_proj = qt_load(m,P("mlp.up_proj.weight"), c->dense_inter, D, dbits); + l->down_proj = qt_load(m,P("mlp.down_proj.weight"), D, c->dense_inter, dbits); + } else { + l->router=ld(m,P("mlp.gate.weight")); + l->router_bias=ld(m,P("mlp.gate.e_score_correction_bias")); + int sI=c->moe_inter*c->n_shared; + l->sh_gate = qt_load(m,P("mlp.shared_experts.gate_proj.weight"), sI, D, dbits); + l->sh_up = qt_load(m,P("mlp.shared_experts.up_proj.weight"), sI, D, dbits); + l->sh_down = qt_load(m,P("mlp.shared_experts.down_proj.weight"), D, sI, dbits); + m->ecache[i]=calloc(cap,sizeof(ESlot)); + m->eroute[i]=calloc(c->topk,sizeof(int)); /* metodo C: ultimo routing del layer */ + m->eusage[i]=calloc(c->n_experts,sizeof(uint32_t)); + } + #undef P + } + /* testa MTP (layer n_layers): presente solo se convertita con --mtp */ + { + /* MTP attiva SOLO se il set e' COMPLETO (i tensori vivono su 3 shard: durante la + * conversione parziale ne esiste solo una parte). MTP=0 la disabilita comunque. */ + const char *req[]={"eh_proj.weight","enorm.weight","hnorm.weight","shared_head.norm.weight", + "input_layernorm.weight","post_attention_layernorm.weight", + "self_attn.q_a_proj.weight","self_attn.q_b_proj.weight","self_attn.kv_a_proj_with_mqa.weight", + "self_attn.kv_b_proj.weight","self_attn.o_proj.weight","mlp.gate.weight", + "mlp.shared_experts.gate_proj.weight","mlp.shared_experts.down_proj.weight", + "mlp.experts.0.gate_proj.weight","mlp.experts.255.down_proj.weight"}; + char mn[256]; m->has_mtp=1; + for(unsigned q=0;qn_layers,req[q]); + if(!st_has(&m->S,mn)){ m->has_mtp=0; break; } + } + if(getenv("MTP") && atoi(getenv("MTP"))==0) m->has_mtp=0; + if(m->has_mtp){ + int i=c->n_layers; Layer *l=&m->mtpL; + #define PM(s) (snprintf(nm,sizeof(nm),"model.layers.%d." s,i),nm) + l->in_ln=ld(m,PM("input_layernorm.weight")); + l->post_ln=ld(m,PM("post_attention_layernorm.weight")); + l->q_a = qt_load(m,PM("self_attn.q_a_proj.weight"), c->q_lora, D, dbits); + l->q_a_ln= ld(m,PM("self_attn.q_a_layernorm.weight")); + l->q_b = qt_load(m,PM("self_attn.q_b_proj.weight"), H*c->qk_head, c->q_lora, dbits); + l->kv_a = qt_load(m,PM("self_attn.kv_a_proj_with_mqa.weight"), c->kv_lora+c->qk_rope, D, dbits); + l->kv_a_ln= ld(m,PM("self_attn.kv_a_layernorm.weight")); + l->kv_b = qt_load(m,PM("self_attn.kv_b_proj.weight"), H*(c->qk_nope+c->v_head), c->kv_lora, dbits); + l->o = qt_load(m,PM("self_attn.o_proj.weight"), D, H*c->v_head, dbits); + l->sparse=1; + l->router=ld(m,PM("mlp.gate.weight")); + l->router_bias=ld(m,PM("mlp.gate.e_score_correction_bias")); + int sI=c->moe_inter*c->n_shared; + l->sh_gate = qt_load(m,PM("mlp.shared_experts.gate_proj.weight"), sI, D, dbits); + l->sh_up = qt_load(m,PM("mlp.shared_experts.up_proj.weight"), sI, D, dbits); + l->sh_down = qt_load(m,PM("mlp.shared_experts.down_proj.weight"), D, sI, dbits); + m->eh_proj = qt_load(m,PM("eh_proj.weight"), D, 2*D, dbits); + m->enorm=ld(m,PM("enorm.weight")); m->hnorm=ld(m,PM("hnorm.weight")); + m->mtp_norm=ld(m,PM("shared_head.norm.weight")); + m->ecache[i]=calloc(cap,sizeof(ESlot)); + m->eroute[i]=calloc(c->topk,sizeof(int)); + m->eusage[i]=calloc(c->n_experts,sizeof(uint32_t)); + m->kv_start[i]=-1; /* KV MTP: parte dalla prima posizione di decode */ + #undef PM + } + } + m->hlast=falloc(D); m->h_all=falloc((int64_t)64*D); + + /* byte della parte DENSA residente (embed+lm_head+attn+mlp densa+shared+norme) */ + int64_t rb=qt_bytes(&m->embed)+qt_bytes(&m->lm_head); + for(int i=0;in_layers;i++){ Layer *l=&m->L[i]; + rb+=qt_bytes(&l->q_a)+qt_bytes(&l->q_b)+qt_bytes(&l->kv_a)+qt_bytes(&l->kv_b)+qt_bytes(&l->o); + if(!l->sparse) rb+=qt_bytes(&l->gate_proj)+qt_bytes(&l->up_proj)+qt_bytes(&l->down_proj); + else rb+=qt_bytes(&l->sh_gate)+qt_bytes(&l->sh_up)+qt_bytes(&l->sh_down); + } + if(m->has_mtp){ Layer *l=&m->mtpL; + rb+=qt_bytes(&l->q_a)+qt_bytes(&l->q_b)+qt_bytes(&l->kv_a)+qt_bytes(&l->kv_b)+qt_bytes(&l->o); + rb+=qt_bytes(&l->sh_gate)+qt_bytes(&l->sh_up)+qt_bytes(&l->sh_down)+qt_bytes(&m->eh_proj); + } + m->resident_bytes=rb; +} + +/* embed: dequantizza la riga del token (scala per-riga) in x[hidden] */ +static void embed_row(Model *m, int tok, float *x){ + int D=m->c.hidden; QT *e=&m->embed; + if(e->fmt==0){ memcpy(x, e->qf+(int64_t)tok*D, D*sizeof(float)); return; } + if(e->fmt==1){ const int8_t *q=e->q8+(int64_t)tok*D; float s=e->s[tok]; + for(int i=0;ifmt==2){ const uint8_t *q=e->q4+(int64_t)tok*((D+1)/2); float s=e->s[tok]; /* int4 */ + for(int i=0;i>1]; x[i]=(float)((int)(byte&0xF)-8)*s; + if(i+1>4)-8)*s; } + return; } + const uint8_t *q=e->q4+(int64_t)tok*((D+3)/4); float s=e->s[tok]; /* int2 */ + for(int i=0;i>2]; int sh=(i&3)*2; x[i]=(float)((int)((byte>>sh)&3)-2)*s; } +} + +/* carica un expert nello slot. Container pre-quantizzato: le 3 matrici sono contigue nel + * file -> UNA pread coalescente da ~19 MB dentro `slab` (+ le scale in fslab); i QT sono + * viste dentro lo slab (zero copie). Fallback per modelli non quantizzati (oracolo tiny). + * THREAD-SAFE su slot distinti (pread posizionale, st_find read-only). */ +static void expert_load(Model *m, int layer, int eid, ESlot *s){ + Cfg *c=&m->c; int I=c->moe_inter, D=c->hidden, b=m->ebits; + char nm[3][288]; const char *suf[3]={"gate_proj","up_proj","down_proj"}; + for(int k=0;k<3;k++) snprintf(nm[k],sizeof(nm[k]),"model.layers.%d.mlp.experts.%d.%s.weight",layer,eid,suf[k]); + char qn[300]; snprintf(qn,sizeof(qn),"%s.qs",nm[0]); + if(!st_has(&m->S,qn)){ /* fallback: tensori pieni, quantizza a runtime */ + qt_from_disk(m,nm[0],I,D,b,g_drop,&s->g); + qt_from_disk(m,nm[1],I,D,b,g_drop,&s->u); + qt_from_disk(m,nm[2],D,I,b,g_drop,&s->d); + s->eid=eid; return; + } + st_tensor *tw[3], *tq[3]; + for(int k=0;k<3;k++){ + tw[k]=st_find(&m->S,nm[k]); + snprintf(qn,sizeof(qn),"%s.qs",nm[k]); tq[k]=st_find(&m->S,qn); + if(!tw[k]||!tq[k]){ fprintf(stderr,"manca %s\n",nm[k]); exit(1); } + } + int64_t wtot=tw[0]->nbytes+tw[1]->nbytes+tw[2]->nbytes; + int64_t ftot=(tq[0]->nbytes+tq[1]->nbytes+tq[2]->nbytes)/4; + /* rialloca se lo slot (riusato tra layer) e' troppo piccolo per QUESTO expert: + * pread oltre la mappatura = short-read o CORRUZIONE silenziosa dei vicini */ + if(!s->slab || wtot+8192 > s->slab_cap){ + free(s->slab); + if(posix_memalign((void**)&s->slab,4096,wtot+8192)){fprintf(stderr,"OOM slab\n");exit(1);} + s->slab_cap=wtot+8192; + } + if(!s->fslab || ftot > s->fslab_cap){ free(s->fslab); s->fslab=falloc(ftot); s->fslab_cap=ftot; } + int ord[3]={0,1,2}; /* ordina per offset nel file */ + for(int a=0;a<3;a++) for(int bb=a+1;bb<3;bb++) if(tw[ord[bb]]->offoff){ int t=ord[a]; ord[a]=ord[bb]; ord[bb]=t; } + int contig = tw[ord[0]]->fd==tw[ord[1]]->fd && tw[ord[1]]->fd==tw[ord[2]]->fd + && tw[ord[0]]->off+tw[ord[0]]->nbytes==tw[ord[1]]->off + && tw[ord[1]]->off+tw[ord[1]]->nbytes==tw[ord[2]]->off; + int64_t pos[3]; int done=0; + if(contig){ + int64_t off0=tw[ord[0]]->off; + int dfd = g_direct ? st_direct_fd(&m->S, tw[ord[0]]->fd) : -1; + if(dfd>=0){ /* O_DIRECT: offset/len allineati a 4K */ + int64_t base=off0 & ~4095LL, need=(off0-base)+wtot; + int64_t len=(need+4095)&~4095LL; + ssize_t r=pread(dfd, s->slab, len, base); + if(r>=need){ + pos[ord[0]]=off0-base; pos[ord[1]]=pos[ord[0]]+tw[ord[0]]->nbytes; + pos[ord[2]]=pos[ord[1]]+tw[ord[1]]->nbytes; done=1; + } + } + if(!done){ /* fallback bufferizzato */ + if(pread(tw[ord[0]]->fd, s->slab, wtot, off0)!=wtot){ perror("pread expert"); exit(1); } + pos[ord[0]]=0; pos[ord[1]]=tw[ord[0]]->nbytes; pos[ord[2]]=tw[ord[0]]->nbytes+tw[ord[1]]->nbytes; done=1; + } + } + if(!done){ /* non contigui: 3 pread bufferizzate */ + int64_t o=0; + for(int a=0;a<3;a++){ int k=ord[a]; + if(pread(tw[k]->fd, s->slab+o, tw[k]->nbytes, tw[k]->off)!=tw[k]->nbytes){ perror("pread expert"); exit(1); } + pos[k]=o; o+=tw[k]->nbytes; } + } + float *fp[3]; int64_t fo=0; /* scale (piccole) */ + for(int k=0;k<3;k++){ + if(pread(tq[k]->fd, (char*)(s->fslab+fo), tq[k]->nbytes, tq[k]->off)!=tq[k]->nbytes){ perror("pread qs"); exit(1); } + fp[k]=s->fslab+fo; fo+=tq[k]->nbytes/4; } + if(g_drop){ /* scarta subito le pagine: evita che la page + * cache in pressione strangoli il throughput */ + posix_fadvise(tw[ord[0]]->fd, tw[ord[0]]->off, wtot, POSIX_FADV_DONTNEED); + for(int k=0;k<3;k++) posix_fadvise(tq[k]->fd, tq[k]->off, tq[k]->nbytes, POSIX_FADV_DONTNEED); + } + QT *qt[3]={&s->g,&s->u,&s->d}; int OO[3]={I,I,D}, II[3]={D,D,I}; + for(int k=0;k<3;k++){ + int64_t nb=tw[k]->nbytes; + int fmt = (nb==(int64_t)OO[k]*II[k])?1 : (nb==(int64_t)OO[k]*((II[k]+1)/2))?2 : 3; + qt[k]->fmt=fmt; qt[k]->O=OO[k]; qt[k]->I=II[k]; qt[k]->qf=NULL; + qt[k]->q8=(int8_t*)(s->slab+pos[k]); qt[k]->q4=s->slab+pos[k]; qt[k]->s=fp[k]; + } + s->eid=eid; +} + +/* prefetch asincrono dei pesi di un expert (e delle sue scale .qs): avvia il readahead + * cosi' le letture sincrone successive trovano la page-cache calda. */ +static void expert_prefetch(Model *m, int layer, int eid){ + char nm[300]; + const char *suf[3]={"gate_proj.weight","up_proj.weight","down_proj.weight"}; + for(int k=0;k<3;k++){ + snprintf(nm,sizeof(nm),"model.layers.%d.mlp.experts.%d.%s",layer,eid,suf[k]); st_prefetch(&m->S,nm); + char qs[320]; snprintf(qs,sizeof(qs),"%s.qs",nm); st_prefetch(&m->S,qs); + } +} + +/* attenzione MLA con KV-cache compressa, su token nuovi x[S,hidden], pos_base = pos del primo */ +static void attention(Model *m, Layer *l, int layer, float *x, int S, int pos_base, float *out){ + Cfg *c=&m->c; int H=c->n_heads, D=c->hidden, qh=c->qk_head, vh=c->v_head; + int kvb_dim=H*(c->qk_nope+vh), Tk=pos_base+S; + double ta0=now_s(); + float *ctx=falloc((int64_t)S*H*vh); + float *Q=falloc((int64_t)S*H*qh); /* query (roped) dei token nuovi */ + float *qresid=falloc(c->q_lora), *comp=falloc(c->kv_lora+c->qk_rope); + /* 1) per ogni token nuovo: query roped + latente normato e k_rot roped -> in cache */ + for(int s=0;sq_a, 1); + rmsnorm(qresid, qresid, l->q_a_ln, c->q_lora, c->eps); + float *qfull=Q+(int64_t)s*H*qh; matmul_qt(qfull, qresid, &l->q_b, 1); + for(int h=0;hqk_nope, pos, c); + matmul_qt(comp, xs, &l->kv_a, 1); + float *Ldst=m->Lc[layer]+(int64_t)pos*c->kv_lora, *Rdst=m->Rc[layer]+(int64_t)pos*c->qk_rope; + memcpy(Ldst, comp, c->kv_lora*sizeof(float)); + rmsnorm(Ldst, Ldst, l->kv_a_ln, c->kv_lora, c->eps); /* latente normato */ + memcpy(Rdst, comp+c->kv_lora, c->qk_rope*sizeof(float)); + rope_interleave(Rdst, pos, c); /* k_rot roped, condiviso fra teste */ + } + /* 2) ricostruzione di k_nope+value per TUTTI i token 0..Tk-1 (un solo matmul su kv_b) */ + double tk0=now_s(); + int stL=m->kv_start[layer]; + float *kvb_all=falloc((int64_t)Tk*kvb_dim); + matmul_qt(kvb_all+(int64_t)stL*kvb_dim, m->Lc[layer]+(int64_t)stL*c->kv_lora, &l->kv_b, Tk-stL); + m->t_kvb += now_s()-tk0; + /* 3) attenzione causale: score = q_pass·k_nope + q_rot·k_rot */ + #pragma omp parallel for collapse(2) schedule(static) + for(int s=0;sqk_nope; + float sc[8192]; + int st0=m->kv_start[layer]; + for(int t=st0;t<=pos;t++){ + const float *kn=kvb_all+(int64_t)t*kvb_dim+(int64_t)h*(c->qk_nope+vh); + const float *kr=m->Rc[layer]+(int64_t)t*c->qk_rope; + float a=0; for(int d=0;dqk_nope;d++) a+=qp[d]*kn[d]; + for(int d=0;dqk_rope;d++) a+=qr[d]*kr[d]; + sc[t-st0]=a*c->attn_scale; + } + softmax(sc,pos+1-st0); + float *cx=ctx+((int64_t)s*H+h)*vh; for(int d=0;dqk_nope+vh)+c->qk_nope; + float a=sc[t-st0]; for(int d=0;do, S); + free(ctx); free(Q); free(qresid); free(comp); free(kvb_all); + m->t_attn += now_s()-ta0; +} + +/* MoE GLM su x[S,hidden] -> out (router sigmoid/noaux_tc, n_group=1, + shared expert). + * BATCH-UNION: per S>1 (prefill, verifica MTP) ogni expert UNICO del batch viene caricato + * una volta sola e moltiplicato per tutte le posizioni che lo usano (pesi letti 1 volta); + * lo shared expert e' un unico matmul a S righe. Per posizione l'accumulo resta + * nell'ordine (routed nel loro ordine di union, poi shared). */ +static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out){ + Cfg *c=&m->c; int D=c->hidden, E=c->n_experts, K=c->topk, I=c->moe_inter; + float *logit=falloc(E), *sig=falloc(E), *choice=falloc(E); + int sI=c->moe_inter*c->n_shared; + /* ---- FASE A: routing di tutte le S posizioni ---- */ + int *idxs=malloc((size_t)S*K*sizeof(int)); float *ws=malloc((size_t)S*K*sizeof(float)); + int *keff=malloc(S*sizeof(int)); + for(int s=0;srouter, 1, D, E); + for(int e=0;erouter_bias[e]; } + int *idx=idxs+(int64_t)s*K; float *w=ws+(int64_t)s*K; + int Ksel = g_topk>0 ? (g_topkbv){bv=choice[e];best=e;} } + idx[kk]=best; w[kk]=sig[best]; + } + int Ke=Ksel; + if(g_topp>0 && g_topp<1.f){ + for(int a=1;a=0 && w[b]=g_topp*tot){ Ke=kk+1; break; } } + } + keff[s]=Ke; m->ereq+=Ke; + for(int kk=0;kkeusage[layer][idx[kk]]++; + if(c->norm_topk){ float sm=0; for(int kk=0;kkrouted_scale; + for(int d=0;denr[layer]=keff[S-1]; for(int kk=0;kkeroute[layer][kk]=idxs[(int64_t)(S-1)*K+kk]; + /* ---- FASE B: union degli expert del batch ---- */ + int *uniq=malloc((size_t)E*sizeof(int)); int nu=0; + { char *seen=calloc(E,1); + for(int s=0;spin[layer]; + for(int z=0;znpin[layer];z++) if(P[z].eid==eid){ m->hits++; use[j]=&P[z]; break; } + if(!use[j]){ ESlot *Sl=m->ecache[layer]; int nn=m->ecn[layer]; + for(int z=0;zhits++; Sl[z].used=++m->eclock; use[j]=&Sl[z]; break; } } + if(!use[j]){ use[j]=&m->ws[nmiss]; missk[nmiss++]=j; m->miss++; } + } + if(nmiss){ double t0=now_s(); + #pragma omp parallel for schedule(dynamic,1) + for(int q=0;qws[q]); + m->t_edisk += now_s()-t0; } + for(int j=0;jg, nr); + matmul_qt(uu, xg, &e->u, nr); + for(int64_t z=0;z<(int64_t)nr*I;z++) gg[z]=siluf(gg[z])*uu[z]; + matmul_qt(hh, gg, &e->d, nr); + for(int r=0;rt_emm += now_s()-t0; + } + { ESlot *Sl=m->ecache[layer]; int *nn=&m->ecn[layer]; /* promozione LRU (swap buffer) */ + int promo = nmissecap ? nmiss : m->ecap; + for(int a=0;aecap) dst=&Sl[(*nn)++]; + else { int lru=0; for(int z=1;z<*nn;z++) if(Sl[z].usedws[q]; m->ws[q]=tmp; dst->used=++m->eclock; } + } + } + /* ---- FASE E: shared expert, un matmul a S righe ---- */ + float *sg=falloc((int64_t)S*sI), *su=falloc((int64_t)S*sI); + matmul_qt(sg, x, &l->sh_gate, S); + matmul_qt(su, x, &l->sh_up, S); + for(int64_t z=0;z<(int64_t)S*sI;z++) sg[z]=siluf(sg[z])*su[z]; + matmul_qt(hh, sg, &l->sh_down, S); + for(int64_t z=0;z<(int64_t)S*D;z++) out[z]+=hh[z]; + free(logit); free(sig); free(choice); free(idxs); free(ws); free(keff); free(uniq); + free(xg); free(gg); free(uu); free(hh); free(rows); free(rw); free(sg); free(su); +} + +static void dense_mlp(Layer *l, float *x, int S, int D, int I, float *out){ + float *g=falloc((int64_t)S*I), *u=falloc((int64_t)S*I); + matmul_qt(g, x, &l->gate_proj, S); + matmul_qt(u, x, &l->up_proj, S); + for(int64_t i=0;i<(int64_t)S*I;i++) g[i]=siluf(g[i])*u[i]; + matmul_qt(out, g, &l->down_proj, S); + free(g); free(u); +} + +/* forward di UN layer (usato dai 78 principali e dal layer MTP) */ +static void layer_forward(Model *m, Layer *l, int li, float *x, int S, int pos_base, float *nrm, float *tmp){ + Cfg *c=&m->c; int D=c->hidden; + if(g_spec && g_prefetch && l->sparse && m->enr[li]>0) + for(int z=0;zenr[li];z++) expert_prefetch(m,li,m->eroute[li][z]); + for(int s=0;sin_ln, D, c->eps); + attention(m,l,li,nrm,S,pos_base,tmp); + for(int64_t j=0;j<(int64_t)S*D;j++) x[j]+=tmp[j]; + for(int s=0;spost_ln, D, c->eps); + if(l->sparse) moe(m,l,li,nrm,S,tmp); else dense_mlp(l,nrm,S,D,c->dense_inter,tmp); + for(int64_t j=0;j<(int64_t)S*D;j++) x[j]+=tmp[j]; +} +static void layers_forward(Model *m, float *x, int S, int pos_base){ + Cfg *c=&m->c; int D=c->hidden; + float *nrm=falloc((int64_t)S*D), *tmp=falloc((int64_t)S*D); + for(int i=0;in_layers;i++) layer_forward(m,&m->L[i],i,x,S,pos_base,nrm,tmp); + free(nrm); free(tmp); +} + +static void kv_alloc(Model *m, int max_t){ + Cfg *c=&m->c; + if(m->Lc){ for(int i=0;in_layers+1;i++){ free(m->Lc[i]); free(m->Rc[i]); } free(m->Lc); free(m->Rc); } + m->max_t=max_t; + int NR=c->n_layers+1; /* riga extra: KV del layer MTP */ + m->Lc=calloc(NR,sizeof(float*)); m->Rc=calloc(NR,sizeof(float*)); + for(int i=0;iLc[i]=falloc((int64_t)max_t*c->kv_lora); + m->Rc[i]=falloc((int64_t)max_t*c->qk_rope); } +} + +static void mtp_absorb(Model *m, const int *next_ids, const float *x, int S, int pos_base); +static float *step(Model *m, const int *ids, int S, int pos_base){ + Cfg *c=&m->c; int D=c->hidden; + float *x=falloc((int64_t)S*D); + for(int s=0;shlast) memcpy(m->hlast, x+(int64_t)(S-1)*D, D*sizeof(float)); + if(m->has_mtp && S>=2 && g_draft>0) mtp_absorb(m, ids+1, x, S-1, pos_base); + float *last=falloc(D); rmsnorm(last, x+(int64_t)(S-1)*D, m->final_norm, D, c->eps); + double th0=now_s(); + float *logit=falloc(c->vocab); matmul_qt(logit,last,&m->lm_head,1); + m->t_head += now_s()-th0; + free(x); free(last); return logit; +} + +/* come step(), ma ritorna i logits di TUTTE le S posizioni [S,vocab] (per la verifica spec) */ +static float *step_all(Model *m, const int *ids, int S, int pos_base){ + Cfg *c=&m->c; int D=c->hidden; + float *x=falloc((int64_t)S*D); + for(int s=0;sh_all) memcpy(m->h_all, x, (int64_t)S*D*sizeof(float)); /* hidden di TUTTE le pos (S<=64) */ + if(m->hlast) memcpy(m->hlast, x+(int64_t)(S-1)*D, D*sizeof(float)); + float *lo=falloc((int64_t)S*c->vocab), *row=falloc(D); + for(int s=0;sfinal_norm, D, c->eps); + matmul_qt(lo+(int64_t)s*c->vocab, row, &m->lm_head, 1); } + free(x); free(row); return lo; +} + +/* METODO E — prompt-lookup: cerca l'occorrenza piu' recente dell'ultimo bigramma nel + * contesto e propone i token che la seguirono. Zero pesi extra, zero costo: e' solo + * un'ipotesi che il modello verifichera'. */ +static int ngram_draft(const int *ids, int len, int G, int *draft){ + if(len<4 || G<1) return 0; + int a=ids[len-2], b=ids[len-1]; + for(int i=len-3;i>=1;i--) + if(ids[i-1]==a && ids[i]==b){ + int n=0; for(int j=i+1;jbv){bv=lo[i];b=i;} return b; +} +static int mtp_draft(Model *m, int next_tok, int kv, int G, int *draft){ + Cfg *c=&m->c; int D=c->hidden, li=c->n_layers; + int p=kv-1; if(p<0||G<1) return 0; + if(m->kv_start[li]<0 || m->kv_start[li]>p) m->kv_start[li]=p; + float *x=falloc(D), *cat=falloc(2*D), *hx=falloc(D), *nrm=falloc(D), *tmp=falloc(D); + float *row=falloc(D), *logit=falloc(c->vocab), *h=falloc(D); + memcpy(h, m->hlast, D*sizeof(float)); + int tok=next_tok, n=0; + int prenorm = getenv("MTP_PRENORM")!=NULL; + for(int g=0; g=m->max_t) break; + embed_row(m, tok, x); + rmsnorm(x, x, m->enorm, D, c->eps); + if(g==0 && !prenorm) rmsnorm(h, h, m->final_norm, D, c->eps); /* h vero: post model.norm */ + rmsnorm(h, h, m->hnorm, D, c->eps); + if(getenv("MTP_SWAP")){ memcpy(cat, h, D*sizeof(float)); memcpy(cat+D, x, D*sizeof(float)); } + else { memcpy(cat, x, D*sizeof(float)); memcpy(cat+D, h, D*sizeof(float)); } + matmul_qt(hx, cat, &m->eh_proj, 1); + double n_eh=0; for(int d=0;d=2; + int t_pre=-1; + if(dbg){ rmsnorm(row, hx, m->mtp_norm, D, c->eps); matmul_qt(logit, row, &m->lm_head, 1); + t_pre=mtp_argmax(logit, c->vocab); } + layer_forward(m, &m->mtpL, li, hx, 1, pos, nrm, tmp); + double n_post=0; for(int d=0;dmtp_norm, D, c->eps); + matmul_qt(logit, row, &m->lm_head, 1); + int t2=mtp_argmax(logit, c->vocab); + if(dbg) fprintf(stderr,"[mtp2] pos=%d in_tok=%d ||eh||=%.1f ||post||=%.1f pre_blk=%d post_blk=%d\n", + pos, tok, sqrt(n_eh), sqrt(n_post), t_pre, t2); + draft[n++]=t2; tok=t2; memcpy(h, hx, D*sizeof(float)); + } + free(x); free(cat); free(hx); free(nrm); free(tmp); free(row); free(logit); free(h); + return n; +} +/* assorbe nella KV della testa MTP le coppie VERIFICATE (emb(token@pos+1), h_vero@pos): + * next_ids[i] = token alla posizione pos_base+i+1; x[i] = hidden VERO a pos_base+i. + * Un solo passaggio batch del layer MTP (il batch-union rende economici gli expert). */ +static void mtp_absorb(Model *m, const int *next_ids, const float *x, int S, int pos_base){ + if(!m->has_mtp || S<1) return; + Cfg *c=&m->c; int D=c->hidden, li=c->n_layers; + if(m->kv_start[li]<0 || m->kv_start[li]>pos_base) m->kv_start[li]=pos_base; + float *hx=falloc((int64_t)S*D), *cat=falloc(2*D), *e=falloc(D), *hn=falloc(D), *hf=falloc(D); + int prenorm = getenv("MTP_PRENORM")!=NULL; + for(int i=0;ienorm,D,c->eps); + if(prenorm) rmsnorm(hn,x+(int64_t)i*D,m->hnorm,D,c->eps); + else { rmsnorm(hf,x+(int64_t)i*D,m->final_norm,D,c->eps); /* vLLM: h POST model.norm */ + rmsnorm(hn,hf,m->hnorm,D,c->eps); } + if(getenv("MTP_SWAP")){ memcpy(cat,hn,D*sizeof(float)); memcpy(cat+D,e,D*sizeof(float)); } + else { memcpy(cat,e,D*sizeof(float)); memcpy(cat+D,hn,D*sizeof(float)); } + matmul_qt(hx+(int64_t)i*D, cat, &m->eh_proj, 1); + } + float *nrm=falloc((int64_t)S*D), *tmp=falloc((int64_t)S*D); + layer_forward(m,&m->mtpL,li,hx,S,pos_base,nrm,tmp); + free(hx); free(cat); free(e); free(hn); free(hf); free(nrm); free(tmp); +} + +static inline int argmax_v(const float *lo, int V){ + int b=0; float bv=lo[0]; for(int i=1;ibv){bv=lo[i];b=i;} return b; +} + +/* stop-set attivo (popolato da run_text/run_serve dal config; vuoto in validazione, + * dove si genera un numero fisso di token da confrontare con l'oracolo) */ +static int g_stop[9], g_nstop=0; +static inline int is_stop(int t){ for(int i=0;in_stop;i++) g_stop[g_nstop++]=c->stop_ids[i]; + if(tok_eos>=0 && !is_stop(tok_eos)) g_stop[g_nstop++]=tok_eos; + fprintf(stderr,"[stop] %d token di stop:",g_nstop); + for(int i=0;i disco e banda RAM ammortizzati su piu' token. + * 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. */ +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; + while(emitted=0 && next==eos) || is_stop(next)) break; + emit(next,ud); all[kv]=next; emitted++; m->n_emit++; + if(emitted>=n_new) break; /* l'ultimo token non serve forwardarlo */ + int g = 0; + if(g_draft>0){ + /* auto-off adattivo: draft che non vengono mai accettati = solo tassa disco */ + if(m->has_mtp && m->mtp_prop>=24 && m->mtp_acc*10 < m->mtp_prop){ + g_draft=0; + fprintf(stderr,"[MTP] acceptance %.0f%% dopo %llu proposte: draft disattivati\n", + 100.0*m->mtp_acc/m->mtp_prop, (unsigned long long)m->mtp_prop); + } + } + if(g_draft>0){ + if(m->has_mtp){ g=mtp_draft(m,next,kv,g_draft,draft); m->mtp_prop+=g; } + else g=ngram_draft(all,kv+1,g_draft,draft); + } + if(g>n_new-emitted) g=n_new-emitted; + if(kv+1+g+1>m->max_t) g=m->max_t-kv-2; + if(g<0) g=0; + int S=1+g; int batch[64]; batch[0]=next; memcpy(batch+1,draft,g*sizeof(int)); + float *lo=step_all(m,batch,S,kv); m->n_fw++; + int k=0; /* verifica: accetta finche' coincide */ + if(g>0 && getenv("MTP_DEBUG")){ int veri=argmax_v(lo,V); + fprintf(stderr,"[mtpdbg] draft0=%d verita=%d %s\n", draft[0], veri, draft[0]==veri?"HIT":"miss"); } + while(k=0 && draft[k]==eos) || is_stop(draft[k])){ done=1; break; } + emit(draft[k],ud); all[kv+1+k]=draft[k]; emitted++; m->n_emit++; k++; + } + if(m->has_mtp) m->mtp_acc+=k; + if(m->has_mtp && k>=1) mtp_absorb(m, all+kv+1, m->h_all, k, kv); /* KV MTP in sync coi verificati */ + /* hlast deve corrispondere all'ultima posizione ACCETTATA (kv+k), non a fine batch */ + if(m->h_all && khlast, m->h_all+(int64_t)k*m->c.hidden, m->c.hidden*sizeof(float)); + kv += 1+k; /* KV oltre kv e' stantia: verra' sovrascritta */ + logit=falloc(V); memcpy(logit, lo+(int64_t)k*V, V*sizeof(float)); free(lo); + } + if(logit) free(logit); + if(kv_out) *kv_out=kv; + return emitted; +} + +/* emit callback: accumula in un array (validazione) */ +typedef struct { int *dst; int n; } EmitStore; +static void emit_store(int t, void *ud){ EmitStore *e=(EmitStore*)ud; e->dst[e->n++]=t; } +/* emit callback: detokenizza e stampa in streaming (chat/run), con heartbeat */ +typedef struct { Tok *T; Model *m; double t0; int count; int quiet; } EmitStream; +static void emit_stream(int t, void *ud){ + EmitStream *e=(EmitStream*)ud; char dec[64]; + int dn=tok_decode(e->T,&t,1,dec,63); dec[dn]=0; fputs(dec,stdout); fflush(stdout); + if(!e->quiet && ++e->count%16==0){ double tt=e->m->hits+e->m->miss; + fprintf(stderr,"\n[t=%d RSS %.2f GB hit %.0f%% %.2f tok/s %.2f tok/fw]\n", e->count, + rss_gb(), tt?100.0*e->m->hits/tt:0.0, e->count/(now_s()-e->t0), + e->m->n_fw?(double)e->m->n_emit/e->m->n_fw:1.0); } +} + +/* teacher-forcing: un solo forward su ids[S], argmax per posizione in pred[S] */ +static void forward_all(Model *m, const int *ids, int S, int *pred){ + Cfg *c=&m->c; int D=c->hidden; + kv_alloc(m,S); + float *x=falloc((int64_t)S*D); + for(int s=0;svocab); + for(int s=0;sfinal_norm, D, c->eps); + matmul_qt(lo, row, &m->lm_head, 1); + int best=0; float bv=lo[0]; for(int i=1;ivocab;i++) if(lo[i]>bv){bv=lo[i];best=i;} + pred[s]=best; + } + free(x); free(lo); +} + +/* log-prob (log-softmax) del token target dato il vettore di logit; *am=1 se e' l'argmax */ +static double logprob_target(const float *lo, int V, int target, int *am){ + float mx=lo[0]; int best=0; for(int i=1;imx){mx=lo[i];best=i;} } + double se=0; for(int i=0;i .. " (T=ctxlen+contlen) + * output: riga " " per richiesta. + * Un solo forward per richiesta (teacher-forcing): niente generazione -> fattibile a bassa velocita'. */ +static void run_score(Model *m, const char *path){ + Cfg *c=&m->c; int D=c->hidden; + FILE *f=fopen(path,"rb"); if(!f){perror(path);exit(1);} + int maxT=1; { char *ln=NULL; size_t cp=0; + while(getline(&ln,&cp,f)>0){ int a,b; if(sscanf(ln,"%d %d",&a,&b)==2 && a+b>maxT) maxT=a+b; } + free(ln); } + kv_alloc(m,maxT); + float *x=falloc((int64_t)maxT*D), *lo=falloc(c->vocab), *row=falloc(D); + int *ids=malloc(maxT*sizeof(int)); + rewind(f); char *ln=NULL; size_t cp=0; int nreq=0; double t0=now_s(); + while(getline(&ln,&cp,f)>0){ + char *p=ln; int ctxlen=strtol(p,&p,10), contlen=strtol(p,&p,10), T=ctxlen+contlen; + if(T<=0||ctxlen<1){ printf("0 0 0\n"); fflush(stdout); continue; } + for(int i=0;ifinal_norm, D, c->eps); + matmul_qt(lo,row,&m->lm_head,1); + int am; lp += logprob_target(lo,c->vocab,ids[pos+1],&am); if(!am) greedy=0; + } + printf("%.6f %d %d\n", lp, contlen, greedy); fflush(stdout); + if(++nreq%5==0) fprintf(stderr,"[score %d req | %.1fs | RSS %.2f GB | hit %.0f%%]\n", + nreq, now_s()-t0, rss_gb(), (m->hits+m->miss)?100.0*m->hits/(m->hits+m->miss):0.0); + } + free(ln); free(ids); free(x); free(lo); free(row); fclose(f); +} + +static void generate(Model *m, const int *prompt, int np, int n_new, int *out){ + kv_alloc(m,np+n_new+g_draft+2); + for(int i=0;ic; char tkp[2048]; snprintf(tkp,sizeof(tkp),"%s/tokenizer.json",snap); + Tok T; tok_load(&T,tkp); + int eos=tok_id_of(&T,"<|endoftext|>"); + stops_arm(&m->c, eos); + int cap=(int)strlen(prompt)+16; int *pids=malloc(cap*sizeof(int)); + int np=tok_encode(&T,prompt,(int)strlen(prompt),pids,cap); + if(np<1){ fprintf(stderr,"prompt vuoto dopo tokenizzazione\n"); return; } + printf("prompt: %d token | genero fino a %d (stop EOS=%d) | draft n-gram=%d\n", np, ngen, eos, g_draft); + fputs(prompt,stdout); fflush(stdout); + kv_alloc(m, np+ngen+g_draft+2); + int *all=malloc((np+ngen+g_draft+2)*sizeof(int)); memcpy(all,pids,np*sizeof(int)); + double t=now_s(); + float *logit=step(m,pids,np,0); + EmitStream es={&T,m,t,0,0}; + int produced=spec_decode(m,all,np,ngen,eos,logit,emit_stream,&es,NULL); + double dt=now_s()-t; + double tot=m->hits+m->miss; + int nsp=0; for(int i=0;in_layers;i++) if(m->L[i].sparse) nsp++; + printf("\n---\n%d token in %.2fs (%.2f tok/s) | hit-rate expert %.1f%% | RSS %.2f GB\n", + produced, dt, produced/dt, tot?100.0*m->hits/tot:0.0, rss_gb()); + printf("expert caricati/token: %.1f (per-layer %.2f su %d; baseline topk=%d) | TOPK=%d TOPP=%.2f\n", + produced?(double)m->ereq/produced:0.0, (produced&&nsp)?(double)m->ereq/produced/nsp:0.0, nsp, c->topk, g_topk, g_topp); + printf("speculazione: %.2f token/forward (%llu fw per %llu tok) | 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, + m->mtp_prop?100.0*m->mtp_acc/m->mtp_prop:0.0, (unsigned long long)m->mtp_acc, (unsigned long long)m->mtp_prop); + double acc=m->t_edisk+m->t_emm+m->t_attn+m->t_head; + printf("PROFILO: expert-disk %.1fs | expert-matmul %.1fs | attention %.1fs (di cui kvb %.1fs) | lm_head %.1fs | altro %.1fs\n", + m->t_edisk, m->t_emm, m->t_attn, m->t_kvb, m->t_head, dt-acc); + free(pids); free(all); +} + +/* modalita' SERVE (per la CLI 'coli'): carica il modello UNA volta, poi CHAT conversazionale. + * KV-cache PERSISTENTE tra i turni: la storia resta in cache, si fa il prefill solo dei + * token NUOVI -> il modello RICORDA la conversazione e non ri-processa il passato (lossless, + * piu' umano, piu' veloce). Template chat GLM con token speciali (CHAT_TEMPLATE=0 -> grezzo). + * Protocollo: "\x01\x01" "READY" "\x01\x01\n" dopo il load; risposta in streaming; "\x01\x01" "END" "\x01\x01\n" a fine turno. + * ":reset" (riga "\x02RESET") azzera la memoria. EOF -> esce. */ +static void run_serve(Model *m, const char *snap){ + char tkp[2048]; snprintf(tkp,sizeof(tkp),"%s/tokenizer.json",snap); + Tok T; tok_load(&T,tkp); + int eos=tok_id_of(&T,"<|endoftext|>"); + stops_arm(&m->c, eos); + int ngen=getenv("NGEN")?atoi(getenv("NGEN")):256; + int maxctx=getenv("CTX")?atoi(getenv("CTX")):4096; + int templ=getenv("CHAT_TEMPLATE")?atoi(getenv("CHAT_TEMPLATE")):1; + kv_alloc(m,maxctx); + int len=0, first=1; /* len = contesto gia' in KV (persiste tra turni) */ + int *hist=malloc(maxctx*sizeof(int)); /* storia token (= contenuto della KV): serve + * al lookup n-gram e resta allineata a len */ + char *line=NULL; size_t cap=0; ssize_t nr; char *buf=malloc(1<<16); + printf("\x01\x01" "READY" "\x01\x01\n"); printf("STAT 0 0.00 0.0 %.2f\n", rss_gb()); fflush(stdout); + while((nr=getline(&line,&cap,stdin))>0){ + 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; + printf("\x01\x01" "END" "\x01\x01\n"); printf("STAT 0 0.00 0.0 %.2f\n", rss_gb()); fflush(stdout); continue; } + if(nr<1){ printf("\x01\x01" "END" "\x01\x01\n"); printf("STAT 0 0.00 0.0 %.2f\n", rss_gb()); fflush(stdout); continue; } + int bl=0; /* costruisce il testo del turno (con template) */ + /* template UFFICIALE GLM-5.2 (chat_template.jinja): niente \n dopo i ruoli, e dopo + * <|assistant|> serve SEMPRE il blocco think — lo DISATTIVA (nothink): + * col template sbagliato il modello farfuglia e non emette mai lo stop. THINK=1 lo abilita. */ + const char *tk = getenv("THINK")&&atoi(getenv("THINK"))? "" : ""; + if(templ){ if(first) bl+=snprintf(buf+bl,(1<<16)-bl,"[gMASK]"); + bl+=snprintf(buf+bl,(1<<16)-bl,"<|user|>%s<|assistant|>%s",line,tk); } + else bl+=snprintf(buf+bl,(1<<16)-bl,"%s",line); + int k=tok_encode(&T,buf,bl,hist+len,maxctx-len); + if(k<1){ printf("\x01\x01" "END" "\x01\x01\n"); printf("STAT 0 0.00 0.0 %.2f\n", rss_gb()); fflush(stdout); continue; } + if(len+k+8+g_draft>=maxctx){ len=0; first=1; /* contesto pieno: azzera e ricomincia */ + bl=0; if(templ){ bl+=snprintf(buf+bl,(1<<16)-bl,"[gMASK]<|user|>%s<|assistant|>%s",line,tk); } + else bl+=snprintf(buf+bl,(1<<16)-bl,"%s",line); + k=tok_encode(&T,buf,bl,hist,maxctx); if(k>maxctx-8-g_draft) k=maxctx-8-g_draft; } + first=0; + int cur=ngen; if(len+k+cur+g_draft+2>=maxctx) cur=maxctx-len-k-g_draft-2; + uint64_t h0=m->hits, ms0=m->miss; double tt0=now_s(); + float *logit=step(m,hist+len,k,len); len+=k; + EmitStream es={&T,m,now_s(),0,1}; + int prod=0; + if(cur>0) prod=spec_decode(m,hist,len,cur,eos,logit,emit_stream,&es,&len); + else free(logit); + double tdt=now_s()-tt0; if(tdt<1e-6) tdt=1e-6; + double dh=(double)(m->hits-h0), dm=(double)(m->miss-ms0); + printf("\n\x01\x01" "END" "\x01\x01\n"); + printf("STAT %d %.2f %.1f %.2f\n", prod, prod/tdt, (dh+dm)>0?100.0*dh/(dh+dm):0.0, rss_gb()); + fflush(stdout); + } + free(line); free(hist); free(buf); +} + +static int *read_arr(jval*o,const char*k,int*n){ jval*a=json_get(o,k); int*r=malloc(a->len*sizeof(int)); + for(int i=0;ilen;i++) r[i]=(int)a->kids[i]->num; *n=a->len; return r; } + +/* byte residenti di un tensore [O,I] al numero di bit dato (specchio di qt_bytes) */ +static int64_t tbytes(int O,int I,int bits){ + if(bits>=16) return (int64_t)O*I*4; + if(bits>=5) return (int64_t)O*I + (int64_t)O*4; + return (int64_t)O*((I+1)/2) + (int64_t)O*4; +} +/* byte VERI di un expert: dal container se pre-quantizzato, altrimenti stima da ebits */ +static int64_t expert_bytes_probe(Model *m, int ebits){ + Cfg *c=&m->c; int64_t eb=0; char nm[256]; + snprintf(nm,sizeof(nm),"model.layers.%d.mlp.experts.0.gate_proj.weight",c->first_dense); + if(st_nbytes(&m->S,nm)>0){ + const char *suf[3]={"gate_proj","up_proj","down_proj"}; + for(int k=0;k<3;k++){ + snprintf(nm,sizeof(nm),"model.layers.%d.mlp.experts.0.%s.weight",c->first_dense,suf[k]); + eb+=st_nbytes(&m->S,nm); + snprintf(nm,sizeof(nm),"model.layers.%d.mlp.experts.0.%s.weight.qs",c->first_dense,suf[k]); + int64_t q=st_nbytes(&m->S,nm); if(q>0) eb+=q; + } + } + if(eb<=0) eb = tbytes(c->moe_inter,c->hidden,ebits)*2 + tbytes(c->hidden,c->moe_inter,ebits); + return eb; +} + +/* scarica su file l'istogramma d'uso degli expert: righe "layer eid count" (per PIN) */ +static void stats_dump(Model *m, const char *path){ + FILE *f=fopen(path,"w"); if(!f){ perror(path); return; } + Cfg *c=&m->c; int64_t tot=0, nz=0; + for(int i=0;in_layers;i++){ if(!m->L[i].sparse) continue; + for(int e=0;en_experts;e++) if(m->eusage[i][e]){ fprintf(f,"%d %d %u\n",i,e,m->eusage[i][e]); tot+=m->eusage[i][e]; nz++; } } + fclose(f); + fprintf(stderr,"[STATS] %lld selezioni su %lld expert distinti -> %s\n",(long long)tot,(long long)nz,path); +} + +/* HOT-STORE ("il redis del colibri'"): carica in RAM, UNA VOLTA e per sempre, i top expert + * per frequenza d'uso misurata (file STATS di un run precedente), entro un budget in GB. + * Ogni hit evita una lettura dal disco lento. */ +static void pin_load(Model *m, const char *statspath, double gb){ + FILE *f=fopen(statspath,"r"); if(!f){ perror(statspath); return; } + typedef struct { int l,e; uint32_t c; } Rec; + Cfg *c=&m->c; int cap=c->n_layers*c->n_experts; + Rec *r=malloc((size_t)cap*sizeof(Rec)); int n=0; + int l,e; uint32_t cnt; + while(n=0&&ln_layers&&e>=0&&en_experts&&m->L[l].sparse) r[n++]=(Rec){l,e,cnt}; + fclose(f); + for(int a=0;ar[best].c) best=b; + Rec t=r[a]; r[a]=r[best]; r[best]=t; + if(a>4095) break; /* bastano i top ~4k */ + } + int64_t eb=expert_bytes_probe(m,m->ebits); + int npin=(int)(gb*1e9/eb); if(npin>n) npin=n; if(npin>4096) npin=4096; + if(npin<1){ free(r); return; } + int *cnt_l=calloc(c->n_layers,sizeof(int)); + for(int a=0;an_layers;i++) if(cnt_l[i]) m->pin[i]=calloc(cnt_l[i],sizeof(ESlot)); + double t0=now_s(); + #pragma omp parallel for schedule(dynamic,1) + for(int a=0;anpin[li]++; + expert_load(m,li,r[a].e,&m->pin[li][slot]); + } + m->resident_bytes += (int64_t)npin*eb; + fprintf(stderr,"[PIN] hot-store: %d expert in RAM (%.1f GB) in %.0fs da %s\n", + npin, npin*eb/1e9, now_s()-t0, statspath); + free(r); free(cnt_l); +} + +static double g_mem_avail_boot=0; /* MemAvailable all'avvio, prima di caricare il modello */ +/* RAM disponibile ADESSO (GB) da /proc/meminfo: e' il tetto vero, non il totale */ +static double mem_available_gb(void){ + FILE *f=fopen("/proc/meminfo","r"); if(!f) return 0; + char ln[256]; double kb=0; + while(fgets(ln,sizeof(ln),f)) if(sscanf(ln,"MemAvailable: %lf",&kb)==1) break; + fclose(f); return kb/1e6; +} + +/* clampa la cache expert a un budget RAM (GB): cap t.c. residente + cache + slack <= budget. + * ram_gb<=0 -> budget AUTO = 88% della RAM disponibile adesso (lascia respiro a OS+wrapper: + * sforare = OOM-kill del kernel a meta' generazione, molto peggio di una cache piu' piccola). */ +static void cap_for_ram(Model *m, double ram_gb, int ebits, int max_ctx){ + Cfg *c=&m->c; int nsp=0; for(int i=0;in_layers;i++) if(m->L[i].sparse) nsp++; + if(m->has_mtp) nsp+=2; /* riga cache MTP: conta ~doppia (expert int8 = 2x int4) */ + int64_t eb=expert_bytes_probe(m,ebits); + int auto_b = ram_gb<=0; + if(auto_b){ ram_gb = g_mem_avail_boot*0.88; /* misurata PRIMA del load: il residente gia' + * allocato viene sottratto sotto, non due volte */ + if(ram_gb<4){ fprintf(stderr,"[RAM] MemAvailable illeggibile/troppo bassa, assumo 8 GB\n"); ram_gb=8; } } + /* slack ONESTO, non forfettario (l'OOM del 2026-07-04 veniva da qui): + * ws[64] slab del working-set (si materializzano TUTTI nel prefill batch-union), + * KV cache a max_ctx, kvb_all della ricostruzione k/v in attention, + * attivazioni+logits+overhead ~1.2 GB */ + double ws_b = 64.0*(double)eb; + double kv_b = (double)(c->n_layers+1)*max_ctx*(c->kv_lora+c->qk_rope)*4.0; + double kvb_b = (double)max_ctx*c->n_heads*(c->qk_nope+c->v_head)*4.0; + double slack = 1.2e9 + ws_b + kv_b + kvb_b; + double avail = ram_gb*1e9 - (double)m->resident_bytes - slack; + int capmax = (avail>0 && nsp>0) ? (int)(avail/((double)nsp*eb)) : 0; + if(capmax<1) capmax=1; + if(capmax < m->ecap){ + fprintf(stderr,"[RAM_GB=%.1f%s] residente %.1f GB + slack %.1f GB (ws %.1f, KV@%d %.1f, kvb %.1f), " + "expert %.1f MB x %d layer -> cap abbassato %d->%d (proiezione picco %.1f GB)\n", + ram_gb, auto_b?" auto":"", m->resident_bytes/1e9, slack/1e9, ws_b/1e9, max_ctx, kv_b/1e9, kvb_b/1e9, + eb/1e6, nsp, m->ecap, capmax, + (m->resident_bytes + (double)capmax*nsp*eb + slack)/1e9); + m->ecap=capmax; + } else { + fprintf(stderr,"[RAM_GB=%.1f%s] cap=%d ok (proiezione picco %.1f GB)\n", ram_gb, auto_b?" auto":"", m->ecap, + (m->resident_bytes + (double)m->ecap*nsp*eb + slack)/1e9); + } +} + +int main(int argc, char **argv){ + /* i thread OMP non devono girare a vuoto mentre il main aspetta il disco */ + if(!getenv("OMP_WAIT_POLICY")) setenv("OMP_WAIT_POLICY","passive",1); + const char *snap=getenv("SNAP"); if(!snap){fprintf(stderr,"SNAP=\n");return 1;} + g_nopack = getenv("NOPACK")?1:0; + g_drop = getenv("DROP")?1:0; + g_prefetch = getenv("PREFETCH")?atoi(getenv("PREFETCH")):0; + g_topk = getenv("TOPK")?atoi(getenv("TOPK")):0; + g_topp = getenv("TOPP")?atof(getenv("TOPP")):0; + g_spec = getenv("SPEC")?atoi(getenv("SPEC")):1; + g_draft = getenv("DRAFT")?atoi(getenv("DRAFT")):-1; /* -1 = auto: 3 se MTP, 0 senza */ + g_direct = getenv("DIRECT")?atoi(getenv("DIRECT")):0; + if(g_draft>63) g_draft=63; /* -1 = auto, risolto dopo model_init */ + int cap = argc>1?atoi(argv[1]):64; + int ebits= argc>2?atoi(argv[2]):8; + int dbits= argc>3?atoi(argv[3]):ebits; + printf("== Motore C GLM (glm_moe_dsa), cache=%d expert/layer | expert@%d-bit densa@%d-bit ==\n", cap, ebits, dbits); + g_mem_avail_boot = mem_available_gb(); + Model m; double t0=now_s(); model_init(&m,snap,cap,ebits,dbits); + if(g_draft<0) g_draft = m.has_mtp ? 3 : 0; + printf("caricato in %.2fs | densa residente: %.2f MB | layers=%d experts=%d | MTP %s (draft=%d)\n", + now_s()-t0, m.resident_bytes/(1024.0*1024.0), m.c.n_layers, m.c.n_experts, + m.has_mtp?"ATTIVA":"assente", g_draft); + if(!strncmp(snap,"/mnt/",5)) + fprintf(stderr,"ATTENZIONE: il modello e' su %s (filesystem 9p/Windows, lento e fadvise inefficace).\n" + " Per RAM e velocita' tienilo su ext4 (es. /home/...).\n", snap); + /* HOT-STORE: PIN= [PIN_GB=g] -> top expert per frequenza fissi in RAM. + * Va PRIMA di cap_for_ram: i pinnati contano nel residente. */ + if(getenv("PIN")) pin_load(&m, getenv("PIN"), getenv("PIN_GB")?atof(getenv("PIN_GB")):10.0); + /* SEMPRE: senza clamp la LRU cresce fino a cap*76 layer = decine di GB -> OOM-kill. + * RAM_GB assente o <=0 = budget automatico da MemAvailable. */ + { int est_ctx = getenv("CTX")?atoi(getenv("CTX")):4096; /* stesso default di run_serve */ + cap_for_ram(&m, getenv("RAM_GB")?atof(getenv("RAM_GB")):0.0, ebits, est_ctx); } + const char *stats=getenv("STATS"); /* STATS= -> istogramma uso expert a fine run */ + + /* modo scoring per benchmark: SCORE= -> log-likelihood per riga */ + if(getenv("SCORE")){ run_score(&m, getenv("SCORE")); if(stats) stats_dump(&m,stats); return 0; } + + /* modo serve persistente per la CLI 'coli': SERVE=1 */ + if(getenv("SERVE")){ run_serve(&m, snap); if(stats) stats_dump(&m,stats); return 0; } + + /* modo testo reale: PROMPT="..." [NGEN=n] -> tokenizza, genera, detokenizza */ + if(getenv("PROMPT")){ + int ngen=getenv("NGEN")?atoi(getenv("NGEN")):64; + run_text(&m, snap, getenv("PROMPT"), ngen); + if(stats) stats_dump(&m,stats); + return 0; + } + + /* altrimenti: validazione contro l'oracolo (ref_glm.json) */ + const char *refpath=getenv("REF")?getenv("REF"):"ref_glm.json"; + FILE *f=fopen(refpath,"rb"); if(!f){perror(refpath);return 1;} + fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); + char *b=malloc(n+1); if(fread(b,1,n,f)!=(size_t)n){} b[n]=0; fclose(f); + char *ar=NULL; jval *ref=json_parse(b,&ar); + int np,nfull; int *prompt=read_arr(ref,"prompt_ids",&np); int *full=read_arr(ref,"full_ids",&nfull); + int n_new=nfull-np; + + if(getenv("TF")){ + int *tf=read_arr(ref,"tf_pred",&(int){0}); + int *pred=malloc(nfull*sizeof(int)); forward_all(&m, full, nfull, pred); + int ok=0; for(int i=0;i [blocco_MB] [n_letture] [threads] [direct 0/1] + * build: gcc -O2 -fopenmp iobench.c -o iobench */ +#define _GNU_SOURCE +#include +#include +#include +#include +#include +#include +#include +#include +#include +static double now(){ struct timespec t; clock_gettime(CLOCK_MONOTONIC,&t); return t.tv_sec+t.tv_nsec*1e-9; } +int main(int argc,char**argv){ + if(argc<2){fprintf(stderr,"uso: %s file [blkMB] [n] [threads] [direct 0/1]\n",argv[0]);return 1;} + long blk=(argc>2?atol(argv[2]):19)*1024*1024; + int n=argc>3?atoi(argv[3]):64; + int nth=argc>4?atoi(argv[4]):8; + int direct=argc>5?atoi(argv[5]):1; + int fd=open(argv[1],O_RDONLY|(direct?O_DIRECT:0)); + if(fd<0 && direct){ fprintf(stderr,"O_DIRECT non disponibile (%s), uso buffered\n",strerror(errno)); + direct=0; fd=open(argv[1],O_RDONLY); } + if(fd<0){perror("open");return 1;} + off_t sz=lseek(fd,0,SEEK_END); + if(sz %.2f GB/s (%.1f ms/blocco effettivi)\n", + direct?"O_DIRECT":"buffered", nth, n, blk/1024/1024, tot/1e9, dt, tot/1e9/dt, dt/n*1000); + close(fd); free(offs); return 0; +} diff --git a/c/json.h b/c/json.h new file mode 100644 index 0000000..3e1c13d --- /dev/null +++ b/c/json.h @@ -0,0 +1,149 @@ +/* Parser JSON minimale, header-only. Serve per: + * - l'header dei file safetensors (un grande oggetto nome->{dtype,shape,data_offsets}) + * - ref.json (per leggere prompt_ids / full_ids) + * Non e' completo (niente unicode \uXXXX, niente notazione esotica) ma copre cio' che serve. */ +#ifndef JSON_H +#define JSON_H +#include +#include +#include +#include + +typedef enum { J_NULL, J_BOOL, J_NUM, J_STR, J_ARR, J_OBJ } jtype; + +typedef struct jval { + jtype t; + double num; /* J_NUM */ + int boolean; /* J_BOOL */ + char *str; /* J_STR (NUL-terminata, dentro l'arena) */ + /* array: figli in [0..len); oggetto: chiavi[] e figli[] in parallelo */ + struct jval **kids; + char **keys; /* solo per J_OBJ */ + int len; +} jval; + +typedef struct { + const char *s; + char *arena; /* buffer per le stringhe smontate */ + size_t acap, aoff; +} jparser; + +static char *j_dup(jparser *p, const char *b, int n) { + /* ogni stringa ha la sua allocazione: un'arena con realloc sposterebbe il + * buffer invalidando i puntatori gia' emessi (use-after-free). */ + (void)p; + char *d = (char *)malloc(n + 1); + memcpy(d, b, n); d[n] = 0; + return d; +} + +static void j_ws(jparser *p) { while (*p->s && isspace((unsigned char)*p->s)) p->s++; } + +static jval *j_new(jtype t) { + jval *v = (jval *)calloc(1, sizeof(jval)); + v->t = t; return v; +} + +static jval *j_parse_val(jparser *p); + +static char *j_parse_str_raw(jparser *p) { + /* assume *p->s == '"' */ + p->s++; + const char *start = p->s; + /* trova la fine gestendo gli escape, poi copia decodificando i casi base */ + char tmp[1 << 16]; int n = 0; + #define J_PUT(ch) do{ if (n < (int)sizeof(tmp)-1) tmp[n++] = (char)(ch); }while(0) + while (*p->s && *p->s != '"') { + char c = *p->s++; + if (c == '\\' && *p->s) { + char e = *p->s++; + switch (e) { + case 'n': c = '\n'; break; case 't': c = '\t'; break; + case 'r': c = '\r'; break; case 'b': c = '\b'; break; + case 'f': c = '\f'; break; case '/': c = '/'; break; + case '\\': c = '\\'; break; case '"': c = '"'; break; + case 'u': { /* \uXXXX -> codepoint UTF-8 (con coppie surrogate) */ + unsigned cp = (unsigned)strtoul((char[]){p->s[0],p->s[1],p->s[2],p->s[3],0}, NULL, 16); + p->s += 4; + if (cp >= 0xD800 && cp <= 0xDBFF && p->s[0]=='\\' && p->s[1]=='u') { + unsigned lo = (unsigned)strtoul((char[]){p->s[2],p->s[3],p->s[4],p->s[5],0}, NULL, 16); + if (lo >= 0xDC00 && lo <= 0xDFFF) { cp = 0x10000 + ((cp-0xD800)<<10) + (lo-0xDC00); p->s += 6; } + } + if (cp < 0x80) { J_PUT(cp); } + else if (cp < 0x800) { J_PUT(0xC0|(cp>>6)); J_PUT(0x80|(cp&0x3F)); } + else if (cp < 0x10000) { J_PUT(0xE0|(cp>>12)); J_PUT(0x80|((cp>>6)&0x3F)); J_PUT(0x80|(cp&0x3F)); } + else { J_PUT(0xF0|(cp>>18)); J_PUT(0x80|((cp>>12)&0x3F)); J_PUT(0x80|((cp>>6)&0x3F)); J_PUT(0x80|(cp&0x3F)); } + continue; + } + default: c = e; break; + } + } + J_PUT(c); + } + #undef J_PUT + if (*p->s == '"') p->s++; + (void)start; + return j_dup(p, tmp, n); +} + +static jval *j_parse_val(jparser *p) { + j_ws(p); + char c = *p->s; + if (c == '"') { jval *v = j_new(J_STR); v->str = j_parse_str_raw(p); return v; } + if (c == '{') { + p->s++; jval *v = j_new(J_OBJ); + int cap = 8; v->keys = malloc(cap * sizeof(char*)); v->kids = malloc(cap * sizeof(jval*)); + j_ws(p); + if (*p->s == '}') { p->s++; return v; } + for (;;) { + j_ws(p); + char *key = j_parse_str_raw(p); + j_ws(p); if (*p->s == ':') p->s++; + jval *val = j_parse_val(p); + if (v->len == cap) { cap *= 2; v->keys = realloc(v->keys, cap*sizeof(char*)); v->kids = realloc(v->kids, cap*sizeof(jval*)); } + v->keys[v->len] = key; v->kids[v->len] = val; v->len++; + j_ws(p); + if (*p->s == ',') { p->s++; continue; } + if (*p->s == '}') { p->s++; break; } + break; + } + return v; + } + if (c == '[') { + p->s++; jval *v = j_new(J_ARR); + int cap = 8; v->kids = malloc(cap * sizeof(jval*)); + j_ws(p); + if (*p->s == ']') { p->s++; return v; } + for (;;) { + jval *val = j_parse_val(p); + if (v->len == cap) { cap *= 2; v->kids = realloc(v->kids, cap*sizeof(jval*)); } + v->kids[v->len++] = val; + j_ws(p); + if (*p->s == ',') { p->s++; continue; } + if (*p->s == ']') { p->s++; break; } + break; + } + return v; + } + if (c == 't') { p->s += 4; jval *v = j_new(J_BOOL); v->boolean = 1; return v; } + if (c == 'f') { p->s += 5; jval *v = j_new(J_BOOL); v->boolean = 0; return v; } + if (c == 'n') { p->s += 4; return j_new(J_NULL); } + /* numero */ + { char *end; double d = strtod(p->s, &end); p->s = end; jval *v = j_new(J_NUM); v->num = d; return v; } +} + +/* API */ +static jval *json_parse(const char *text, char **arena_out) { + jparser p = { text, NULL, 0, 0 }; + jval *v = j_parse_val(&p); + if (arena_out) *arena_out = p.arena; else free(p.arena); + return v; +} + +static jval *json_get(jval *o, const char *key) { + if (!o || o->t != J_OBJ) return NULL; + for (int i = 0; i < o->len; i++) if (strcmp(o->keys[i], key) == 0) return o->kids[i]; + return NULL; +} + +#endif diff --git a/c/make_glm_oracle.py b/c/make_glm_oracle.py new file mode 100644 index 0000000..9eb6780 --- /dev/null +++ b/c/make_glm_oracle.py @@ -0,0 +1,79 @@ +"""Costruisce un GLM-5.2 (glm_moe_dsa) MINUSCOLO a pesi random come ORACOLO. +Architettura vera (MLA + DSA indexer + router sigmoid/noaux_tc + shared expert), +dimensioni minuscole. Salva pesi+config in c/glm_tiny/ e un riferimento greedy in +c/ref_glm.json. seq corta (<= index_topk) cosi' il DSA seleziona tutte le key e +l'attenzione coincide con la MLA densa: il motore C puo' validare senza implementare +l'indexer sparso.""" +import json, torch +from transformers import GlmMoeDsaConfig, GlmMoeDsaForCausalLM + +torch.manual_seed(1234) + +cfg = GlmMoeDsaConfig( + vocab_size=256, + hidden_size=128, + intermediate_size=64, # MLP densa (primi 3 layer) + moe_intermediate_size=32, # expert + num_hidden_layers=5, # 3 densi + 2 sparse + first_k_dense_replace=3, + num_attention_heads=4, + num_key_value_heads=4, + n_routed_experts=8, + num_experts_per_tok=2, + n_shared_experts=1, + q_lora_rank=64, + kv_lora_rank=32, + qk_nope_head_dim=24, + qk_rope_head_dim=8, # pari -> interleave ok; head_dim diventa 8 + v_head_dim=32, + index_topk=4096, # >> seq_len -> DSA seleziona tutto (no-op) + index_head_dim=16, + index_n_heads=2, + n_group=1, + topk_group=1, + norm_topk_prob=True, + routed_scaling_factor=2.5, + rope_parameters={"rope_type": "default", "rope_theta": 10000.0}, + tie_word_embeddings=False, + rms_norm_eps=1e-5, + attention_bias=False, + max_position_embeddings=4096, +) +cfg._attn_implementation = "eager" + +model = GlmMoeDsaForCausalLM(cfg).eval() +# rende i pesi non banali (default init e' molto piccolo): scala router/bias per topk vario +with torch.no_grad(): + for n, p in model.named_parameters(): + if p.dim() >= 2: + p.normal_(0, 0.05) + # bias di correzione del router: valori distinti cosi' la selezione e' sensata + for i, layer in enumerate(model.model.layers): + if hasattr(layer.mlp, "gate"): + layer.mlp.gate.e_score_correction_bias.copy_( + torch.linspace(-0.1, 0.1, cfg.n_routed_experts)) + +print("=== tensori dello state_dict (nomi per il loader C) ===") +for n, p in model.state_dict().items(): + print(f" {n:60s} {tuple(p.shape)}") + +prompt = [3, 14, 159, 26, 53, 58, 200, 11, 77, 240, 5, 99] # token id arbitrari, seq corta +ids = torch.tensor([prompt]) +with torch.no_grad(): + out = model.generate(ids, max_new_tokens=20, do_sample=False, use_cache=True) +full = out[0].tolist() +print("\nprompt:", prompt) +print("full :", full) + +# teacher-forcing: un singolo forward su tutta la sequenza -> argmax per posizione. +# Per il greedy vale tf_pred[i] == full[i+1] per i >= len(prompt)-1; serve a validare +# il PREFILL del motore C separandolo dal decode. +with torch.no_grad(): + lg = model(torch.tensor([full]), use_cache=False).logits[0] # [seq, vocab] +tf_pred = lg.argmax(-1).tolist() +print("tf_pred:", tf_pred) + +model.save_pretrained("glm_tiny", safe_serialization=True) +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("\nsalvato: glm_tiny/ (pesi+config) e ref_glm.json") diff --git a/c/olmoe.c b/c/olmoe.c new file mode 100644 index 0000000..3d2a07d --- /dev/null +++ b/c/olmoe.c @@ -0,0 +1,390 @@ +/* Motore di inferenza OLMoE in C puro, con EXPERT-STREAMING dal disco. + * Porting del motore Python (engine.py). Obiettivo Stadio A: produrre gli STESSI + * token id del riferimento (ref.json) -> valida il core prima di scalare a GLM-5.2. + * + * Densa (embed, attn, router, norme, lm_head) residente in RAM (float32). + * Expert letti dal disco on-demand via pread+fadvise(DONTNEED), cache LRU per-layer. + * Matmul multi-thread con OpenMP (niente BLAS). + */ +#define _GNU_SOURCE +#include +#include +#include +#include +#include +#include +#include "st.h" + +/* ---------- config ---------- */ +typedef struct { + int hidden, n_layers, n_heads, n_kv_heads, head_dim; + int n_experts, topk, inter, vocab; + float theta, eps; int norm_topk; +} Cfg; + +/* ---------- pesi densi per-layer ---------- */ +typedef struct { + float *in_ln, *post_ln, *q, *k, *v, *o, *qn, *kn, *gate; +} Layer; + +/* ---------- cache LRU degli expert (pesi QUANTIZZATI) ---------- + * Ogni weight [out,in] tenuto come int8 (per-riga) + scala float per riga. + * Cosi' la RAM-cache scende da 4 byte/param (f32) a 1 byte/param: e' il + * meccanismo che fa stare GLM-5.2 nei 15 GB. dequant-on-use nel matmul. */ +typedef struct { int eid; int8_t *g, *u, *d; float *gs, *us, *ds; uint64_t used; } Slot; +typedef struct { Slot *slots; int n, cap; } LCache; + +typedef struct { + Cfg c; + shards S; + int quant_bits; /* bit di quantizzazione degli expert (2..8); 16 = f32 */ + float *embed, *lm_head, *final_norm; + Layer *L; + LCache *cache; /* [n_layers] */ + uint64_t clock, hits, miss; + /* kv-cache per-layer: K,V come [H * maxT * head_dim] */ + float **K, **V; int kv_len, max_t; + double dense_load_s; +} Model; + +/* ---------- utility ---------- */ +static double now_s(void) { struct timespec t; clock_gettime(CLOCK_MONOTONIC, &t); return t.tv_sec + t.tv_nsec*1e-9; } +static double rss_gb(void) { struct rusage r; getrusage(RUSAGE_SELF, &r); return r.ru_maxrss / (1024.0*1024.0); } +static float *falloc(int64_t n) { float *p = malloc(n*sizeof(float)); if(!p){fprintf(stderr,"OOM %ld\n",(long)n);exit(1);} return p; } + +/* y[S,O] = x[S,I] @ W^T, W e' [O,I] row-major */ +static void matmul(float *y, const float *x, const float *W, int S, int I, int O) { + #pragma omp parallel for schedule(static) + for (int o = 0; o < O; o++) { + const float *w = W + (int64_t)o * I; + for (int s = 0; s < S; s++) { + const float *xs = x + (int64_t)s * I; + float acc = 0.f; + for (int i = 0; i < I; i++) acc += xs[i] * w[i]; + y[(int64_t)s * O + o] = acc; + } + } +} + +/* y[1,O] = x[1,I] @ W^T con W quantizzato: q[O,I] int8 + scala per riga. + * W[o,i] ~= q[o,i]*scale[o] -> y[o] = scale[o] * sum_i x[i]*q[o,i]. */ +static void matmul_q(float *y, const float *x, const int8_t *q, const float *scale, int I, int O) { + #pragma omp parallel for schedule(static) + for (int o = 0; o < O; o++) { + const int8_t *w = q + (int64_t)o * I; + float acc = 0.f; + for (int i = 0; i < I; i++) acc += x[i] * (float)w[i]; + y[o] = acc * scale[o]; + } +} + +/* quantizza un weight f32 [O,I] -> int8 q[O,I] + scala[O], simmetrica per riga. + * Replica quant_dequant() del Python: scale = amax(|w|, riga)/qmax, q = round(w/scale). */ +static void quantize_rows(const float *w, int8_t *q, float *scale, int O, int I, int bits) { + int qmax = (1 << (bits - 1)) - 1; /* 8->127, 4->7, 2->1 */ + #pragma omp parallel for schedule(static) + for (int o = 0; o < O; o++) { + const float *wr = w + (int64_t)o * I; + float amax = 0.f; for (int i = 0; i < I; i++) { float a = fabsf(wr[i]); if (a > amax) amax = a; } + float s = amax / qmax; if (s < 1e-8f) s = 1e-8f; + scale[o] = s; + int8_t *qr = q + (int64_t)o * I; + for (int i = 0; i < I; i++) { + int v = (int)lrintf(wr[i] / s); + if (v > qmax) v = qmax; + if (v < -qmax-1) v = -qmax-1; + qr[i] = (int8_t)v; + } + } +} + +/* rmsnorm su una riga di lunghezza D, in-place su out (out puo' essere == x) */ +static void rmsnorm_row(float *out, const float *x, const float *w, int D, float eps) { + double ms = 0; for (int i = 0; i < D; i++) ms += (double)x[i]*x[i]; + float r = 1.f / sqrtf((float)(ms / D) + eps); + for (int i = 0; i < D; i++) out[i] = x[i] * r * w[i]; +} + +static void softmax_row(float *x, int n) { + float m = -1e30f; for (int i = 0; i < n; i++) if (x[i] > m) m = x[i]; + float s = 0; for (int i = 0; i < n; i++) { x[i] = expf(x[i]-m); s += x[i]; } + for (int i = 0; i < n; i++) x[i] /= s; +} + +/* ---------- caricamento ---------- */ +static void load_cfg(Cfg *c, const char *snap) { + char path[2048]; snprintf(path, sizeof(path), "%s/config.json", snap); + FILE *f = fopen(path, "rb"); if(!f){perror(path);exit(1);} + fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); + char *buf = malloc(n+1); if(fread(buf,1,n,f)!=(size_t)n){} buf[n]=0; fclose(f); + char *arena=NULL; jval *r = json_parse(buf, &arena); + c->hidden = (int)json_get(r,"hidden_size")->num; + c->n_layers = (int)json_get(r,"num_hidden_layers")->num; + c->n_heads = (int)json_get(r,"num_attention_heads")->num; + c->n_kv_heads= (int)json_get(r,"num_key_value_heads")->num; + c->n_experts = (int)json_get(r,"num_experts")->num; + c->topk = (int)json_get(r,"num_experts_per_tok")->num; + c->inter = (int)json_get(r,"intermediate_size")->num; + c->vocab = (int)json_get(r,"vocab_size")->num; + c->head_dim = c->hidden / c->n_heads; + jval *th = json_get(r,"rope_theta"); c->theta = th ? (float)th->num : 10000.f; + jval *ep = json_get(r,"rms_norm_eps"); c->eps = ep ? (float)ep->num : 1e-5f; + jval *nt = json_get(r,"norm_topk_prob"); c->norm_topk = (nt && nt->t==J_BOOL) ? nt->boolean : 0; + free(buf); free(arena); +} + +static float *load_t(Model *m, const char *name) { + int64_t n = st_numel(&m->S, name); + if (n < 0) { fprintf(stderr, "manca %s\n", name); exit(1); } + float *p = falloc(n); + st_read_f32(&m->S, name, p, 0); /* densa: niente DONTNEED, resta residente */ + return p; +} + +static void model_init(Model *m, const char *snap, int cap, int bits) { + memset(m, 0, sizeof(*m)); + m->quant_bits = bits; + load_cfg(&m->c, snap); + st_init(&m->S, snap); + Cfg *c = &m->c; + double t0 = now_s(); + m->embed = load_t(m, "model.embed_tokens.weight"); + m->lm_head = load_t(m, "lm_head.weight"); + m->final_norm = load_t(m, "model.norm.weight"); + m->L = calloc(c->n_layers, sizeof(Layer)); + char nm[256]; + for (int i = 0; i < c->n_layers; i++) { + Layer *l = &m->L[i]; + #define LD(field, suffix) snprintf(nm,sizeof(nm),"model.layers.%d." suffix,i); l->field = load_t(m,nm) + LD(in_ln, "input_layernorm.weight"); + LD(post_ln,"post_attention_layernorm.weight"); + LD(q, "self_attn.q_proj.weight"); LD(k, "self_attn.k_proj.weight"); + LD(v, "self_attn.v_proj.weight"); LD(o, "self_attn.o_proj.weight"); + LD(qn,"self_attn.q_norm.weight"); LD(kn,"self_attn.k_norm.weight"); + LD(gate, "mlp.gate.weight"); + #undef LD + } + m->cache = calloc(c->n_layers, sizeof(LCache)); + for (int i = 0; i < c->n_layers; i++) { m->cache[i].cap = cap; m->cache[i].slots = calloc(cap, sizeof(Slot)); } + m->dense_load_s = now_s() - t0; +} + +/* legge un weight dal disco (streaming) e lo quantizza in q[O,I]+scale[O] */ +static void load_expert_w(Model *m, const char *name, int8_t *q, float *scale, int O, int I, float *tmp) { + st_read_f32(&m->S, name, tmp, 1); /* pread + fadvise DONTNEED */ + quantize_rows(tmp, q, scale, O, I, m->quant_bits); +} + +/* ---------- cache expert: ritorna i pesi quantizzati (q+scale) da cache o disco ---------- */ +static void expert_get(Model *m, int layer, int eid, Slot **out) { + LCache *lc = &m->cache[layer]; + for (int i = 0; i < lc->n; i++) if (lc->slots[i].eid == eid) { + m->hits++; lc->slots[i].used = ++m->clock; *out = &lc->slots[i]; return; + } + m->miss++; + Cfg *c = &m->c; + int64_t ng = (int64_t)c->inter * c->hidden, nd = (int64_t)c->hidden * c->inter; + Slot *s; + if (lc->n < lc->cap) { + s = &lc->slots[lc->n++]; + s->g = malloc(ng); s->u = malloc(ng); s->d = malloc(nd); + s->gs = falloc(c->inter); s->us = falloc(c->inter); s->ds = falloc(c->hidden); + } else { int lru = 0; for (int i = 1; i < lc->n; i++) if (lc->slots[i].used < lc->slots[lru].used) lru = i; s = &lc->slots[lru]; } + float *tmp = falloc(ng > nd ? ng : nd); + char nm[256]; + snprintf(nm,sizeof(nm),"model.layers.%d.mlp.experts.%d.gate_proj.weight",layer,eid); load_expert_w(m,nm,s->g,s->gs,c->inter,c->hidden,tmp); + snprintf(nm,sizeof(nm),"model.layers.%d.mlp.experts.%d.up_proj.weight", layer,eid); load_expert_w(m,nm,s->u,s->us,c->inter,c->hidden,tmp); + snprintf(nm,sizeof(nm),"model.layers.%d.mlp.experts.%d.down_proj.weight",layer,eid); load_expert_w(m,nm,s->d,s->ds,c->hidden,c->inter,tmp); + free(tmp); + s->eid = eid; s->used = ++m->clock; + *out = s; +} + +/* ---------- RoPE su un vettore di una testa (head_dim) a posizione assoluta pos ---------- */ +static void rope_head(float *x, int pos, const Cfg *c) { + int h = c->head_dim / 2; + for (int j = 0; j < h; j++) { + float inv = powf(c->theta, -2.0f * j / c->head_dim); + float ang = pos * inv, cs = cosf(ang), sn = sinf(ang); + float a = x[j], b = x[j+h]; + x[j] = a*cs - b*sn; + x[j+h] = b*cs + a*sn; + } +} + +/* attenzione sui token nuovi x[S,hidden]; pos_base = posizione assoluta del primo token nuovo */ +static void attention(Model *m, Layer *l, int layer, float *x, int S, int pos_base, float *out) { + Cfg *c = &m->c; int H = c->n_heads, hd = c->head_dim, D = c->hidden; + float *q = falloc((int64_t)S*D), *k = falloc((int64_t)S*D), *vv = falloc((int64_t)S*D); + matmul(q, x, l->q, S, D, D); + matmul(k, x, l->k, S, D, D); + matmul(vv, x, l->v, S, D, D); + /* qk-norm sull'intero vettore hidden, poi RoPE per testa */ + for (int s = 0; s < S; s++) { + rmsnorm_row(q + (int64_t)s*D, q + (int64_t)s*D, l->qn, D, c->eps); + rmsnorm_row(k + (int64_t)s*D, k + (int64_t)s*D, l->kn, D, c->eps); + int pos = pos_base + s; + for (int hh = 0; hh < H; hh++) { rope_head(q + (int64_t)s*D + hh*hd, pos, c); rope_head(k + (int64_t)s*D + hh*hd, pos, c); } + } + /* scrive k,v nella kv-cache alle posizioni pos_base..pos_base+S-1 */ + for (int s = 0; s < S; s++) for (int hh = 0; hh < H; hh++) { + int t = pos_base + s; + memcpy(m->K[layer] + ((int64_t)hh*m->max_t + t)*hd, k + (int64_t)s*D + hh*hd, hd*sizeof(float)); + memcpy(m->V[layer] + ((int64_t)hh*m->max_t + t)*hd, vv + (int64_t)s*D + hh*hd, hd*sizeof(float)); + } + int Tk = pos_base + S; /* numero di key totali disponibili */ + float scale = 1.f / sqrtf((float)hd); + float *ctx = falloc((int64_t)S*D); + #pragma omp parallel for collapse(2) schedule(static) + for (int hh = 0; hh < H; hh++) { + for (int s = 0; s < S; s++) { + int qpos = pos_base + s; + const float *qv = q + (int64_t)s*D + hh*hd; + float sc[4096]; + for (int t = 0; t <= qpos; t++) { /* causale: t <= qpos */ + const float *kv = m->K[layer] + ((int64_t)hh*m->max_t + t)*hd; + float acc = 0; for (int dd = 0; dd < hd; dd++) acc += qv[dd]*kv[dd]; + sc[t] = acc * scale; + } + softmax_row(sc, qpos+1); + float *cx = ctx + (int64_t)s*D + hh*hd; + for (int dd = 0; dd < hd; dd++) cx[dd] = 0; + for (int t = 0; t <= qpos; t++) { + const float *vrow = m->V[layer] + ((int64_t)hh*m->max_t + t)*hd; + float a = sc[t]; + for (int dd = 0; dd < hd; dd++) cx[dd] += a * vrow[dd]; + } + } + } + (void)Tk; + matmul(out, ctx, l->o, S, D, D); + free(q); free(k); free(vv); free(ctx); +} + +/* MoE sui token x[S,hidden] -> out[S,hidden] */ +static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out) { + Cfg *c = &m->c; int D = c->hidden, E = c->n_experts, K = c->topk, I = c->inter; + float *logits = falloc((int64_t)S*E); + matmul(logits, x, l->gate, S, D, E); + memset(out, 0, (int64_t)S*D*sizeof(float)); + float *g = falloc(I), *u = falloc(I), *hh = falloc(D); + for (int s = 0; s < S; s++) { + float *pr = logits + (int64_t)s*E; + softmax_row(pr, E); + /* top-K indici (selezione parziale) */ + int idx[64]; float val[64]; + for (int kk = 0; kk < K; kk++) { + int best = -1; float bv = -1e30f; + for (int e = 0; e < E; e++) { + int taken = 0; for (int j = 0; j < kk; j++) if (idx[j]==e){taken=1;break;} + if (!taken && pr[e] > bv) { bv = pr[e]; best = e; } + } + idx[kk] = best; val[kk] = bv; + } + if (c->norm_topk) { float sm=0; for(int kk=0;kkg, e->gs, D, I); /* gate_proj [I,D] */ + matmul_q(u, xs, e->u, e->us, D, I); /* up_proj [I,D] */ + for (int i = 0; i < I; i++) { float gv = g[i]; g[i] = (gv / (1.f + expf(-gv))) * u[i]; } + matmul_q(hh, g, e->d, e->ds, I, D); /* down_proj [D,I] */ + float w = val[kk]; + float *os = out + (int64_t)s*D; + for (int d = 0; d < D; d++) os[d] += w * hh[d]; + } + } + free(logits); free(g); free(u); free(hh); +} + +/* un passo: token nuovi ids[S] a posizione pos_base. Ritorna logits dell'ultimo token (malloc'd). */ +static float *step(Model *m, const int *ids, int S, int pos_base) { + Cfg *c = &m->c; int D = c->hidden; + float *x = falloc((int64_t)S*D); + for (int s = 0; s < S; s++) memcpy(x + (int64_t)s*D, m->embed + (int64_t)ids[s]*D, D*sizeof(float)); + float *nrm = falloc((int64_t)S*D), *tmp = falloc((int64_t)S*D); + for (int i = 0; i < c->n_layers; i++) { + Layer *l = &m->L[i]; + for (int s = 0; s < S; s++) rmsnorm_row(nrm + (int64_t)s*D, x + (int64_t)s*D, l->in_ln, D, c->eps); + attention(m, l, i, nrm, S, pos_base, tmp); + for (int64_t j = 0; j < (int64_t)S*D; j++) x[j] += tmp[j]; + for (int s = 0; s < S; s++) rmsnorm_row(nrm + (int64_t)s*D, x + (int64_t)s*D, l->post_ln, D, c->eps); + moe(m, l, i, nrm, S, tmp); + for (int64_t j = 0; j < (int64_t)S*D; j++) x[j] += tmp[j]; + } + m->kv_len = pos_base + S; + /* solo l'ultimo token -> logits */ + float *last = falloc(D); + rmsnorm_row(last, x + (int64_t)(S-1)*D, m->final_norm, D, c->eps); + float *logit = falloc(c->vocab); + matmul(logit, last, m->lm_head, 1, D, c->vocab); + free(x); free(nrm); free(tmp); free(last); + return logit; +} + +/* generazione greedy. prompt[np] -> riempie out[np+n_new] */ +static void generate(Model *m, const int *prompt, int np, int n_new, int *out) { + Cfg *c = &m->c; + m->max_t = np + n_new; + m->K = calloc(c->n_layers, sizeof(float*)); m->V = calloc(c->n_layers, sizeof(float*)); + for (int i = 0; i < c->n_layers; i++) { + m->K[i] = falloc((int64_t)c->n_heads * m->max_t * c->head_dim); + m->V[i] = falloc((int64_t)c->n_heads * m->max_t * c->head_dim); + } + for (int i = 0; i < np; i++) out[i] = prompt[i]; + float *logit = step(m, prompt, np, 0); /* PREFILL */ + int len = np; + for (int s = 0; s < n_new; s++) { + int best = 0; float bv = logit[0]; + for (int i = 1; i < c->vocab; i++) if (logit[i] > bv) { bv = logit[i]; best = i; } + free(logit); + out[len++] = best; + if (s == n_new - 1) break; + int one = best; + logit = step(m, &one, 1, len - 1); /* DECODE */ + } +} + +/* ---------- lettura ref.json ---------- */ +static int *read_int_array(jval *o, const char *key, int *n_out) { + jval *a = json_get(o, key); + int *r = malloc(a->len * sizeof(int)); + for (int i = 0; i < a->len; i++) r[i] = (int)a->kids[i]->num; + *n_out = a->len; return r; +} + +int main(int argc, char **argv) { + const char *snap = getenv("SNAP"); + if (!snap) { fprintf(stderr, "imposta SNAP=\n"); return 1; } + int cap = argc > 1 ? atoi(argv[1]) : 16; + int bits = argc > 2 ? atoi(argv[2]) : 8; + const char *refpath = argc > 3 ? argv[3] : "ref.json"; + + FILE *f = fopen(refpath, "rb"); if(!f){perror(refpath);return 1;} + fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); + char *buf=malloc(n+1); if(fread(buf,1,n,f)!=(size_t)n){} buf[n]=0; fclose(f); + char *arena=NULL; jval *ref = json_parse(buf, &arena); + int np, nfull; int *prompt = read_int_array(ref,"prompt_ids",&np); int *full = read_int_array(ref,"full_ids",&nfull); + int n_new = nfull - np; + + printf("== Motore C streaming, cache = %d expert/layer, expert @ %d-bit ==\n", cap, bits); + Model m; model_init(&m, snap, cap, bits); + printf("densa caricata in %.1fs | RSS dopo load densa: %.2f GB\n", m.dense_load_s, rss_gb()); + + int *out = malloc((np + n_new) * sizeof(int)); + double t = now_s(); + generate(&m, prompt, np, n_new, out); + double dt = now_s() - t; + + int match = 0; + printf("\nRiferimento: "); for (int i=np;i/dev/null 2>&1; do + echo "[1/4] attendo lo spostamento su ext4... ($(du -sh "$DIR" 2>/dev/null | cut -f1))"; sleep 20 +done +echo "[1/4] spostamento completato: $(du -sh "$DIR" | cut -f1), shard $(ls "$DIR"/*.safetensors 2>/dev/null | wc -l)" + +# 2) riprende+completa la conversione (ripartibile: salta gli shard gia' fatti) +echo "[2/4] conversione (riprende da dove era): output -> $DIR" +python3 convert_fp8_to_int4.py --repo "$REPO" --outdir "$DIR" --ebits 4 --io-bits 8 + +# 3) il motore richiede tokenizer.json + config.json nella dir del modello +for f in config.json tokenizer.json; do + [ -f "$DIR/$f" ] || { echo "ERRORE: manca $DIR/$f"; exit 1; } +done +echo "[3/4] compilo il motore"; make -s glm + +# 4) generazione reale, con auto-cap dal budget RAM e heartbeat RSS su stderr +echo "[4/4] genero (RAM_GB=$RAM_GB, NGEN=$NGEN)"; echo "------" +SNAP="$DIR" RAM_GB="$RAM_GB" PROMPT="$PROMPT" NGEN="$NGEN" ./glm 64 diff --git a/c/setup.sh b/c/setup.sh new file mode 100644 index 0000000..41e6d08 --- /dev/null +++ b/c/setup.sh @@ -0,0 +1,35 @@ +#!/usr/bin/env bash +# colibrì — installazione su una macchina nuova (Linux x86-64). +# Compila il motore e fa un self-test. Il MODELLO (~372 GB int4) va copiato a parte +# o rigenerato con: coli convert --model +set -e +cd "$(dirname "$0")" +echo "🐦 colibrì — setup" + +# 1) dipendenze +command -v gcc >/dev/null || { echo "manca gcc (es: sudo apt install build-essential)"; exit 1; } +command -v make >/dev/null || { echo "manca make"; exit 1; } +echo " gcc: $(gcc -dumpversion) · $(nproc) core" +echo -n " OpenMP: "; echo 'int main(){return 0;}' | gcc -fopenmp -xc - -o /tmp/_omp 2>/dev/null && echo ok || { echo "manca (libgomp)"; exit 1; } + +# 2) build: nativa (veloce, per QUESTA macchina). Per un binario da distribuire: make portable +echo " compilo (ARCH=${ARCH:-native})…" +make -s glm ARCH="${ARCH:-native}" + +# 3) self-test sull'oracolo tiny, se presente +if [ -d glm_tiny ] && [ -f ref_glm.json ]; then + r=$(SNAP=./glm_tiny TF=1 ./glm 64 16 16 2>/dev/null | grep -oE "[0-9]+/[0-9]+ posizioni" || true) + echo " self-test motore: ${r:-?} (atteso 32/32)" +fi + +# 4) info macchina (la velocità dipende da QUESTI due numeri, non dalla GPU) +ram=$(awk '/MemTotal/{printf "%.0f", $2/1e6}' /proc/meminfo 2>/dev/null || echo "?") +echo " RAM: ${ram} GB (più RAM = più expert in cache = più veloce)" +echo +echo "pronto. Prossimi passi:" +echo " ./coli build # (gia' fatto)" +echo " ./coli convert --model /percorso/NVMe/glm52_i4 # genera il modello int4 (ore)" +echo " ./coli info --model /percorso/NVMe/glm52_i4" +echo " ./coli chat --model /percorso/NVMe/glm52_i4 --ram " +echo +echo "IMPORTANTE: tieni il modello su disco VELOCE (NVMe/ext4), MAI su /mnt/c o rete." diff --git a/c/st.h b/c/st.h new file mode 100644 index 0000000..f608481 --- /dev/null +++ b/c/st.h @@ -0,0 +1,226 @@ +/* Indicizzazione e lettura on-demand di tensori da piu' file safetensors. + * Equivale a Shards in engine.py, ma: + * - legge con pread (niente mmap) + posix_fadvise(DONTNEED) -> le pagine NON + * restano residenti nel processo. E' la correzione del bug di RSS: cosi' la + * RAM di picco resta densa+cache, non l'intero modello. (vedi memoria mmap-rss-bug) + * - converte sempre in float32 in uscita (BF16/F16/F32 supportati). */ +#ifndef ST_H +#define ST_H +#define _GNU_SOURCE +#include +#include +#include +#include +#include +#include +#include +#include "json.h" + +typedef struct { + char *name; + int fd; + int64_t off; /* offset assoluto del dato dentro al file */ + int64_t nbytes; + int dtype; /* 0=BF16 1=F16 2=F32 */ + int64_t numel; +} st_tensor; + +typedef struct { + st_tensor *t; + int n, cap; + int fds[512]; + int dfds[512]; /* gemelli O_DIRECT (aperti pigramente): -2 = non ancora provato */ + char *paths[512]; + int nfd; + int *hidx; /* hash map nome->indice (open addressing): con ~120k tensori + * (GLM: 256 expert x 78 layer x 3 x 2) la scansione lineare + * costava decine di secondi/token (misurato sul primo run reale) */ + int hcap; +} shards; +#define ST_MAX_SHARDS 512 + +static uint64_t st_hash(const char *s){ + uint64_t h=1469598103934665603ULL; + while(*s){ h^=(unsigned char)*s++; h*=1099511628211ULL; } + return h; +} + +static int st_dtype_code(const char *s) { + if (!strcmp(s, "BF16")) return 0; + if (!strcmp(s, "F16")) return 1; + if (!strcmp(s, "F32")) return 2; + if (!strcmp(s, "U8")) return 3; /* dati quantizzati (int4 packed / int8) */ + if (!strcmp(s, "I8")) return 3; + fprintf(stderr, "dtype non gestito: %s\n", s); exit(1); +} + +static inline float bf16_to_f32(uint16_t h) { + uint32_t u = (uint32_t)h << 16; float f; memcpy(&f, &u, 4); return f; +} +static inline float f16_to_f32(uint16_t h) { + uint32_t sign = (uint32_t)(h & 0x8000) << 16; + uint32_t exp = (h >> 10) & 0x1F; + uint32_t man = h & 0x3FF; + uint32_t u; + if (exp == 0) { /* subnormale o zero */ + if (man == 0) u = sign; + else { exp = 127 - 15 + 1; while (!(man & 0x400)) { man <<= 1; exp--; } man &= 0x3FF; u = sign | (exp << 23) | (man << 13); } + } else if (exp == 0x1F) { /* inf/nan */ + u = sign | 0x7F800000 | (man << 13); + } else { + u = sign | ((exp - 15 + 127) << 23) | (man << 13); + } + float f; memcpy(&f, &u, 4); return f; +} + +static int st_open_fd(shards *S, const char *path) { + for (int i = 0; i < S->nfd; i++) if (!strcmp(S->paths[i], path)) return S->fds[i]; + int fd = open(path, O_RDONLY); + if (fd < 0) { perror(path); exit(1); } + S->paths[S->nfd] = strdup(path); S->fds[S->nfd] = fd; + S->dfds[S->nfd] = open(path, O_RDONLY | O_DIRECT); /* eager: lookup poi thread-safe */ + S->nfd++; + return fd; +} + +/* fd gemello O_DIRECT dello stesso file (bypassa la page cache: il buffered read su + * ext4-in-VHDX si strozza a ~0.8 GB/s, O_DIRECT arriva a 2.3+; misurato). -1 se non disponibile. */ +static int st_direct_fd(shards *S, int fd) { + for (int i = 0; i < S->nfd; i++) if (S->fds[i] == fd) return S->dfds[i]; + return -1; +} + +/* indicizza tutti i model-*.safetensors in snap_dir */ +static void st_init(shards *S, const char *snap_dir) { + memset(S, 0, sizeof(*S)); + S->cap = 4096; S->t = calloc(S->cap, sizeof(st_tensor)); + /* raccoglie ordinatamente i nomi dei file shard */ + static char files[ST_MAX_SHARDS][1024]; int nf = 0; + DIR *d = opendir(snap_dir); struct dirent *e; + if (!d) { perror(snap_dir); exit(1); } + while ((e = readdir(d))) { + const char *dot = strrchr(e->d_name, '.'); + if (dot && !strcmp(dot, ".safetensors")) { /* model.safetensors o model-0000N-of-... */ + if (nf >= ST_MAX_SHARDS) { fprintf(stderr, "troppi shard (>%d): alza ST_MAX_SHARDS\n", ST_MAX_SHARDS); exit(1); } + snprintf(files[nf++], 1024, "%s/%s", snap_dir, e->d_name); + } + } + closedir(d); + for (int a = 0; a < nf; a++) for (int b = a+1; b < nf; b++) + if (strcmp(files[a], files[b]) > 0) { char tmp[1024]; strcpy(tmp, files[a]); strcpy(files[a], files[b]); strcpy(files[b], tmp); } + + for (int fi = 0; fi < nf; fi++) { + int fd = st_open_fd(S, files[fi]); + uint64_t hlen; + if (pread(fd, &hlen, 8, 0) != 8) { perror("pread hlen"); exit(1); } + char *hdr = malloc(hlen + 1); + if (pread(fd, hdr, hlen, 8) != (ssize_t)hlen) { perror("pread hdr"); exit(1); } + hdr[hlen] = 0; + int64_t data_start = 8 + (int64_t)hlen; + char *arena = NULL; + jval *root = json_parse(hdr, &arena); + for (int i = 0; i < root->len; i++) { + const char *name = root->keys[i]; + if (!strcmp(name, "__metadata__")) continue; + jval *m = root->kids[i]; + jval *dt = json_get(m, "dtype"); + jval *off = json_get(m, "data_offsets"); + jval *shp = json_get(m, "shape"); + int64_t a0 = (int64_t)off->kids[0]->num, b0 = (int64_t)off->kids[1]->num; + int64_t numel = 1; for (int k = 0; k < shp->len; k++) numel *= (int64_t)shp->kids[k]->num; + if (S->n == S->cap) { S->cap *= 2; S->t = realloc(S->t, S->cap*sizeof(st_tensor)); } + st_tensor *t = &S->t[S->n++]; + t->name = strdup(name); t->fd = fd; t->off = data_start + a0; + t->nbytes = b0 - a0; t->dtype = st_dtype_code(dt->str); t->numel = numel; + } + free(arena); /* i jval restano leakati: ok, una tantum all'avvio */ + free(hdr); + } + /* indice hash costruito a fine indicizzazione (gli indici restano validi dopo i realloc) */ + S->hcap = 1; while (S->hcap < S->n * 2) S->hcap <<= 1; + S->hidx = malloc(S->hcap * sizeof(int)); + for (int i = 0; i < S->hcap; i++) S->hidx[i] = -1; + for (int i = 0; i < S->n; i++) { + uint64_t h = st_hash(S->t[i].name) & (S->hcap - 1); + while (S->hidx[h] >= 0) h = (h + 1) & (S->hcap - 1); + S->hidx[h] = i; + } +} + +static st_tensor *st_find(shards *S, const char *name) { + if (S->hidx) { + uint64_t h = st_hash(name) & (S->hcap - 1); + while (S->hidx[h] >= 0) { + st_tensor *t = &S->t[S->hidx[h]]; + if (!strcmp(t->name, name)) return t; + h = (h + 1) & (S->hcap - 1); + } + return NULL; + } + for (int i = 0; i < S->n; i++) if (!strcmp(S->t[i].name, name)) return &S->t[i]; + return NULL; +} +static int st_has(shards *S, const char *name) { return st_find(S, name) != NULL; } + +/* prefetch ASINCRONO: dice al kernel di iniziare a leggere le pagine del tensore in + * background (readahead). Serve a sovrapporre l'I/O degli expert col calcolo: si + * prefetcha tutto il set di expert di un layer, poi le pread sincrone trovano la cache + * gia' calda. No-op se il tensore non esiste (es. il primo .qs prima della lettura). */ +static void st_prefetch(shards *S, const char *name) { + st_tensor *t = st_find(S, name); + if (t) posix_fadvise(t->fd, t->off, t->nbytes, POSIX_FADV_WILLNEED); +} + +/* legge un tensore in un buffer float32 fornito dal chiamante (numel float). + * drop=1 -> consiglia al kernel di scartare le pagine (per gli expert in streaming). */ +static int64_t st_read_f32(shards *S, const char *name, float *out, int drop) { + st_tensor *t = st_find(S, name); + if (!t) { fprintf(stderr, "tensore mancante: %s\n", name); exit(1); } + void *raw = malloc(t->nbytes); + if (pread(t->fd, raw, t->nbytes, t->off) != t->nbytes) { perror("pread data"); exit(1); } + if (t->dtype == 2) { + memcpy(out, raw, t->nbytes); + } else if (t->dtype == 0) { + uint16_t *p = (uint16_t *)raw; for (int64_t i = 0; i < t->numel; i++) out[i] = bf16_to_f32(p[i]); + } else { + uint16_t *p = (uint16_t *)raw; for (int64_t i = 0; i < t->numel; i++) out[i] = f16_to_f32(p[i]); + } + free(raw); + if (drop) posix_fadvise(t->fd, t->off, t->nbytes, POSIX_FADV_DONTNEED); + return t->numel; +} + +static int64_t st_numel(shards *S, const char *name) { + st_tensor *t = st_find(S, name); return t ? t->numel : -1; +} +static int64_t st_nbytes(shards *S, const char *name) { + st_tensor *t = st_find(S, name); return t ? t->nbytes : -1; +} + +/* legge i byte GREZZI di un tensore (nessuna conversione di dtype): per i pesi gia' + * quantizzati int4/int8 del nostro container (dtype U8). drop=1 -> fadvise DONTNEED. */ +static void st_read_raw(shards *S, const char *name, void *out, int drop) { + st_tensor *t = st_find(S, name); + if (!t) { fprintf(stderr, "tensore mancante: %s\n", name); exit(1); } + if (pread(t->fd, out, t->nbytes, t->off) != t->nbytes) { perror("pread raw"); exit(1); } + if (drop) posix_fadvise(t->fd, t->off, t->nbytes, POSIX_FADV_DONTNEED); +} + +/* legge una FETTA di un tensore: n_elems a partire dall'elemento elem_off. + * Serve per gli expert fusi di GLM (un tensore = blocco [E, ...]): si legge il + * solo expert richiesto via pread del sotto-range, niente lettura dell'intero blocco. */ +static void st_read_slice_f32(shards *S, const char *name, int64_t elem_off, int64_t n_elems, float *out, int drop) { + st_tensor *t = st_find(S, name); + if (!t) { fprintf(stderr, "tensore mancante: %s\n", name); exit(1); } + int esz = (t->dtype == 2) ? 4 : 2; + int64_t boff = t->off + elem_off * esz, nb = n_elems * esz; + void *raw = malloc(nb); + if (pread(t->fd, raw, nb, boff) != nb) { perror("pread slice"); exit(1); } + if (t->dtype == 2) memcpy(out, raw, nb); + else if (t->dtype == 0) { uint16_t *p = raw; for (int64_t i = 0; i < n_elems; i++) out[i] = bf16_to_f32(p[i]); } + else { uint16_t *p = raw; for (int64_t i = 0; i < n_elems; i++) out[i] = f16_to_f32(p[i]); } + free(raw); + if (drop) posix_fadvise(t->fd, boff, nb, POSIX_FADV_DONTNEED); +} + +#endif diff --git a/c/supervisor.sh b/c/supervisor.sh new file mode 100644 index 0000000..144aa32 --- /dev/null +++ b/c/supervisor.sh @@ -0,0 +1,55 @@ +#!/usr/bin/env bash +# Supervisore della conversione GLM-5.2 — a prova di rete WSL che si blocca. +# - tiene SEMPRE vivo un (solo) convertitore +# - se un download resta FERMO >180s (connessione zombie), lo ammazza e lo rilancia: +# hf_hub riprende il .incomplete dal punto esatto, non si perde nulla +# - esce da solo quando tutti i 141 shard sono fatti +# uso: nohup ./supervisor.sh > supervisor.log 2>&1 & +set -u +DIR=/home/vincenzo/glm52_i4 +CODE=/mnt/c/Users/User/Desktop/moe-stream/c +TOTAL=141 +STALL_S=180 # secondi senza crescita del download -> riavvio +CONVLOG=/tmp/convert_supervised.log + +exec 9>"$DIR/.supervisor.lock" +flock -n 9 || { echo "supervisore gia' attivo, esco"; exit 1; } + +log(){ echo "[$(date +%H:%M:%S)] $*"; } + +start_conv(){ + cd "$CODE" + nohup python3 convert_fp8_to_int4.py --repo zai-org/GLM-5.2-FP8 \ + --outdir "$DIR" --ebits 4 --io-bits 8 >> "$CONVLOG" 2>&1 & + log "convertitore avviato (PID $!)" +} + +last_size=-1; stall=0 +while :; do + done_n=$(ls "$DIR"/out-*.safetensors 2>/dev/null | wc -l) + if [ "$done_n" -ge "$TOTAL" ]; then log "FATTO: $done_n/$TOTAL shard. Esco."; pkill -f convert_fp8 2>/dev/null; exit 0; fi + + if ! pgrep -f convert_fp8 >/dev/null; then + log "convertitore non attivo ($done_n/$TOTAL): lo avvio" + start_conv; last_size=-1; stall=0; sleep 20; continue + fi + + inc=$(find "$DIR/_inflight" -name "*.incomplete" 2>/dev/null | head -1) + if [ -n "$inc" ]; then + size=$(stat -c%s "$inc" 2>/dev/null || echo 0) + if [ "$size" = "$last_size" ]; then + stall=$((stall+30)) + if [ "$stall" -ge "$STALL_S" ]; then + log "download FERMO da ${stall}s a $((size/1000000)) MB ($done_n/$TOTAL): riavvio il convertitore" + pkill -f convert_fp8; sleep 5 + start_conv; last_size=-1; stall=0 + fi + else + [ "$last_size" -ge 0 ] && [ "$stall" -ge 60 ] && log "download ripreso ($((size/1000000)) MB)" + last_size=$size; stall=0 + fi + else + last_size=-1; stall=0 # niente .incomplete = sta convertendo/salvando: tutto ok + fi + sleep 30 +done diff --git a/c/tok.h b/c/tok.h new file mode 100644 index 0000000..e5ec6f8 --- /dev/null +++ b/c/tok.h @@ -0,0 +1,278 @@ +/* Tokenizer GLM-5.2 in C puro (byte-level BPE stile cl100k / tiktoken). + * Replica fedele di tokenizer.json: + * - model.type = BPE, ignore_merges=true, byte_fallback=false + * - pre_tokenizer: regex Split (pattern cl100k) + ByteLevel(add_prefix_space=false) + * - merges con rank = ordine nella lista; \p{L}/\p{N}/\s da tok_unicode.h + * - added_tokens (speciali e non) trattati come atomici in encode/decode + * API: + * tok_load(&T, "tokenizer.json"); + * int n = tok_encode(&T, text, len, out_ids, max); + * int m = tok_decode(&T, ids, n, out_buf, max); + */ +#ifndef TOK_H +#define TOK_H +#define _GNU_SOURCE +#include +#include +#include +#include +#include +#include "json.h" +#include "tok_unicode.h" + +/* ---------- hash map (chiavi binarie con lunghezza) ---------- */ +typedef struct { const char *k; int klen; int v; int used; } ment; +typedef struct { ment *e; int cap; } hmap; +static uint64_t tk_fnv(const char *s, int n){ uint64_t h=1469598103934665603ULL; + for(int i=0;icap=cap; m->e=(ment*)calloc(cap,sizeof(ment)); } +static void hm_put(hmap *m, const char *k, int klen, int v){ + uint64_t h=tk_fnv(k,klen)&(m->cap-1); + while(m->e[h].used){ if(m->e[h].klen==klen && !memcmp(m->e[h].k,k,klen)){ m->e[h].v=v; return; } h=(h+1)&(m->cap-1); } + m->e[h].k=k; m->e[h].klen=klen; m->e[h].v=v; m->e[h].used=1; +} +static int hm_get(hmap *m, const char *k, int klen){ + uint64_t h=tk_fnv(k,klen)&(m->cap-1); + while(m->e[h].used){ if(m->e[h].klen==klen && !memcmp(m->e[h].k,k,klen)) return m->e[h].v; h=(h+1)&(m->cap-1); } + return -1; +} + +typedef struct { char *str; int len; int id; } Special; +typedef struct { + hmap vocab; /* stringa byte-level -> id */ + hmap merges; /* "left\0right" -> rank */ + char **id2str; int *id_added; int n_ids; /* id -> stringa; id_added=1 se added-token (output letterale) */ + Special *sp; int nsp; /* added tokens, ordinati per lunghezza decrescente */ + uint32_t byte2cp[256]; int byte2cp_len[256]; char byte2str[256][3]; + int16_t cp2byte[1024]; +} Tok; + +/* ---------- UTF-8 ---------- */ +static int u8_next(const unsigned char *s, int len, int i, uint32_t *cp){ + unsigned char c=s[i]; + if(c<0x80){ *cp=c; return 1; } + if((c>>5)==0x6 && i+1>4)==0xE && i+2>3)==0x1E && i+3>6); o[1]=0x80|(cp&0x3F); return 2; } + if(cp<0x10000){ o[0]=0xE0|(cp>>12); o[1]=0x80|((cp>>6)&0x3F); o[2]=0x80|(cp&0x3F); return 3; } + o[0]=0xF0|(cp>>18); o[1]=0x80|((cp>>12)&0x3F); o[2]=0x80|((cp>>6)&0x3F); o[3]=0x80|(cp&0x3F); return 4; +} + +/* ---------- mappa byte<->unicode di GPT-2/ByteLevel ---------- */ +static void tk_build_bytemap(Tok *T){ + for(int i=0;i<1024;i++) T->cp2byte[i]=-1; + int isdir[256]; memset(isdir,0,sizeof(isdir)); + for(int b=33;b<=126;b++) isdir[b]=1; + for(int b=161;b<=172;b++) isdir[b]=1; + for(int b=174;b<=255;b++) isdir[b]=1; + int n=0; + for(int b=0;b<256;b++){ + uint32_t cp = isdir[b] ? (uint32_t)b : (uint32_t)(256+n); + if(!isdir[b]) n++; + T->byte2cp[b]=cp; + T->byte2cp_len[b]=u8_put(T->byte2str[b], cp); + if(cp<1024) T->cp2byte[cp]=b; + } +} + +/* ---------- caricamento tokenizer.json ---------- */ +static char *tk_read_file(const char *path, long *out_n){ + FILE *f=fopen(path,"rb"); if(!f){ perror(path); exit(1); } + fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); + char *b=malloc(n+1); if(fread(b,1,n,f)!=(size_t)n){} b[n]=0; fclose(f); if(out_n)*out_n=n; return b; +} +static int cmp_sp_len(const void *a, const void *b){ return ((const Special*)b)->len - ((const Special*)a)->len; } + +static void tok_load(Tok *T, const char *path){ + memset(T,0,sizeof(*T)); + tk_build_bytemap(T); + long fn; char *buf=tk_read_file(path,&fn); + char *arena=NULL; jval *root=json_parse(buf,&arena); + jval *model=json_get(root,"model"); + jval *vocab=json_get(model,"vocab"); + jval *merges=json_get(model,"merges"); + jval *added=json_get(root,"added_tokens"); + if(!vocab||!merges){ fprintf(stderr,"tokenizer.json: manca model.vocab/merges\n"); exit(1); } + + /* id massimo per dimensionare id2str */ + int maxid=0; + for(int i=0;ilen;i++){ int id=(int)vocab->kids[i]->num; if(id>maxid)maxid=id; } + if(added) for(int i=0;ilen;i++){ int id=(int)json_get(added->kids[i],"id")->num; if(id>maxid)maxid=id; } + T->n_ids=maxid+1; + T->id2str=calloc(T->n_ids,sizeof(char*)); + T->id_added=calloc(T->n_ids,sizeof(int)); + + /* vocab: stringa -> id (capacita' potenza di 2, ~2-3x) */ + int vc=1; while(vc < vocab->len*2) vc<<=1; + hm_init(&T->vocab, vc); + for(int i=0;ilen;i++){ + const char *k=vocab->keys[i]; int id=(int)vocab->kids[i]->num; + hm_put(&T->vocab, k, (int)strlen(k), id); + T->id2str[id]=(char*)k; + } + /* merges: "left\0right" -> rank=i */ + int mc=1; while(mc < merges->len*2) mc<<=1; + hm_init(&T->merges, mc); + for(int i=0;ilen;i++){ + jval *pr=merges->kids[i]; + const char *l=pr->kids[0]->str, *r=pr->kids[1]->str; + int ll=(int)strlen(l), rl=(int)strlen(r); + char *key=malloc(ll+1+rl); memcpy(key,l,ll); key[ll]=0; memcpy(key+ll+1,r,rl); + hm_put(&T->merges, key, ll+1+rl, i); + } + /* added tokens (speciali e non): atomici, output letterale */ + if(added){ + T->nsp=added->len; T->sp=calloc(T->nsp,sizeof(Special)); + for(int i=0;ilen;i++){ + jval *a=added->kids[i]; + char *content=json_get(a,"content")->str; int id=(int)json_get(a,"id")->num; + T->sp[i].str=content; T->sp[i].len=(int)strlen(content); T->sp[i].id=id; + T->id2str[id]=content; T->id_added[id]=1; + } + qsort(T->sp,T->nsp,sizeof(Special),cmp_sp_len); /* match piu' lungo per primo */ + } + /* arena/buf restano allocati: le stringhe (j_dup) sono malloc indipendenti e ci servono vive */ + (void)arena; +} + +/* ---------- BPE su un pezzo: byte grezzi [a,b) -> id appesi a out ---------- */ +static void bpe_piece(Tok *T, const unsigned char *p, int a, int b, int *out, int *no, int max){ + int nb=b-a; + /* stringa byte-level (concatenazione di byte2str): <=2 byte per byte di input */ + char *s=malloc(2*nb+1); int sl=0; + for(int i=a;ibyte2str[bb],T->byte2cp_len[bb]); sl+=T->byte2cp_len[bb]; } + s[sl]=0; + /* ignore_merges: se l'intero pezzo e' un token, emettilo diretto */ + int whole=hm_get(&T->vocab,s,sl); + if(whole>=0){ if(*nomerges,kbuf,ll+1+rl); + if(rk>=0 && rkvocab,s+soff[i],slen[i]); + if(id>=0 && *no='A'&&(c)<='Z')?((c)+32):(c)) + int i=0; + while(i=0){ + if(is_L(cp[j])){ while(j run di whitespace fino all'ultimo newline contiguo */ + { + int r=i; while(ri){ int last=-1; for(int j=i;j=0){ i=last+1; bpe_piece(T,p,off[start],off[i],out,no,max); continue; } + /* 6) \s+(?!\S): se seguito da non-spazio lascia l'ultimo ws, altrimenti prendi tutto */ + int end = (r id (split sugli added token, poi pretok+BPE) ---------- */ +static int tok_encode(Tok *T, const char *text, int len, int *out, int max){ + const unsigned char *p=(const unsigned char*)text; int no=0; int i=0; + while(i= i (match piu' lungo) */ + int hitpos=-1, hitlen=0, hitid=-1; + for(int j=i;jnsp;k++){ + int sl=T->sp[k].len; + if(sl>0 && j+sl<=len && !memcmp(p+j,T->sp[k].str,sl)){ hitpos=j; hitlen=sl; hitid=T->sp[k].id; break; } + } + } + int chunk_end = (hitpos<0) ? len : hitpos; + if(chunk_end>i) pretok_chunk(T,p,i,chunk_end,out,&no,max); + if(hitpos<0) break; + if(no"); -1 se assente */ +static int tok_id_of(Tok *T, const char *content){ + for(int i=0;insp;i++) if(!strcmp(T->sp[i].str,content)) return T->sp[i].id; + return -1; +} + +/* ---------- decode: id -> testo (byte-level inverso; added token letterali) ---------- */ +static int tok_decode(Tok *T, const int *ids, int n, char *out, int max){ + int o=0; + for(int i=0;i=T->n_ids||!T->id2str[id]) continue; + const char *s=T->id2str[id]; + if(T->id_added[id]){ int l=(int)strlen(s); for(int j=0;jcp2byte[c]>=0 && ocp2byte[c]; } + } + if(o (legge righe "TEXT\tID,ID,.." da stdin) */ +#define _GNU_SOURCE +#include "tok.h" + +int main(int argc, char **argv){ + if(argc<2){ fprintf(stderr,"uso: %s tokenizer.json < casi\n",argv[0]); return 1; } + Tok T; + tok_load(&T, argv[1]); + fprintf(stderr,"caricato: vocab_ids=%d specials=%d\n", T.n_ids, T.nsp); + char *line=NULL; size_t cap=0; ssize_t nr; + int pass=0, tot=0, dpass=0; + while((nr=getline(&line,&cap,stdin))>=0){ + if(nr>0 && line[nr-1]=='\n'){ line[--nr]=0; } + if(nr==0) continue; + char *tab=strchr(line,'\t'); if(!tab) continue; + *tab=0; const char *text=line; const char *idstr=tab+1; + /* il testo puo' contenere \n e \t codificati come \\n \\t */ + char tbuf[4096]; int tn=0; + for(const char *q=text; *q && tn<4095; q++){ + if(q[0]=='\\' && q[1]=='n'){ tbuf[tn++]='\n'; q++; } + else if(q[0]=='\\' && q[1]=='t'){ tbuf[tn++]='\t'; q++; } + else if(q[0]=='\\' && q[1]=='r'){ tbuf[tn++]='\r'; q++; } + else if(q[0]=='\\' && q[1]=='\\'){ tbuf[tn++]='\\'; q++; } + else tbuf[tn++]=*q; + } + tbuf[tn]=0; + int exp[4096], ne=0; + for(const char *q=idstr; *q; ){ while(*q==','||*q==' ')q++; if(!*q)break; exp[ne++]=atoi(q); while(*q&&*q!=',')q++; } + int got[4096]; int ng=tok_encode(&T,tbuf,tn,got,4096); + int ok = (ng==ne); for(int i=0;i + +static const uint32_t uni_L[][2] = { + {0x41,0x5A},{0x61,0x7A},{0xAA,0xAA},{0xB5,0xB5},{0xBA,0xBA},{0xC0,0xD6}, + {0xD8,0xF6},{0xF8,0x2C1},{0x2C6,0x2D1},{0x2E0,0x2E4},{0x2EC,0x2EC},{0x2EE,0x2EE}, + {0x370,0x374},{0x376,0x377},{0x37A,0x37D},{0x37F,0x37F},{0x386,0x386},{0x388,0x38A}, + {0x38C,0x38C},{0x38E,0x3A1},{0x3A3,0x3F5},{0x3F7,0x481},{0x48A,0x52F},{0x531,0x556}, + {0x559,0x559},{0x560,0x588},{0x5D0,0x5EA},{0x5EF,0x5F2},{0x620,0x64A},{0x66E,0x66F}, + {0x671,0x6D3},{0x6D5,0x6D5},{0x6E5,0x6E6},{0x6EE,0x6EF},{0x6FA,0x6FC},{0x6FF,0x6FF}, + {0x710,0x710},{0x712,0x72F},{0x74D,0x7A5},{0x7B1,0x7B1},{0x7CA,0x7EA},{0x7F4,0x7F5}, + {0x7FA,0x7FA},{0x800,0x815},{0x81A,0x81A},{0x824,0x824},{0x828,0x828},{0x840,0x858}, + {0x860,0x86A},{0x870,0x887},{0x889,0x88E},{0x8A0,0x8C9},{0x904,0x939},{0x93D,0x93D}, + {0x950,0x950},{0x958,0x961},{0x971,0x980},{0x985,0x98C},{0x98F,0x990},{0x993,0x9A8}, + {0x9AA,0x9B0},{0x9B2,0x9B2},{0x9B6,0x9B9},{0x9BD,0x9BD},{0x9CE,0x9CE},{0x9DC,0x9DD}, + {0x9DF,0x9E1},{0x9F0,0x9F1},{0x9FC,0x9FC},{0xA05,0xA0A},{0xA0F,0xA10},{0xA13,0xA28}, + {0xA2A,0xA30},{0xA32,0xA33},{0xA35,0xA36},{0xA38,0xA39},{0xA59,0xA5C},{0xA5E,0xA5E}, + {0xA72,0xA74},{0xA85,0xA8D},{0xA8F,0xA91},{0xA93,0xAA8},{0xAAA,0xAB0},{0xAB2,0xAB3}, + {0xAB5,0xAB9},{0xABD,0xABD},{0xAD0,0xAD0},{0xAE0,0xAE1},{0xAF9,0xAF9},{0xB05,0xB0C}, + {0xB0F,0xB10},{0xB13,0xB28},{0xB2A,0xB30},{0xB32,0xB33},{0xB35,0xB39},{0xB3D,0xB3D}, + {0xB5C,0xB5D},{0xB5F,0xB61},{0xB71,0xB71},{0xB83,0xB83},{0xB85,0xB8A},{0xB8E,0xB90}, + {0xB92,0xB95},{0xB99,0xB9A},{0xB9C,0xB9C},{0xB9E,0xB9F},{0xBA3,0xBA4},{0xBA8,0xBAA}, + {0xBAE,0xBB9},{0xBD0,0xBD0},{0xC05,0xC0C},{0xC0E,0xC10},{0xC12,0xC28},{0xC2A,0xC39}, + {0xC3D,0xC3D},{0xC58,0xC5A},{0xC5D,0xC5D},{0xC60,0xC61},{0xC80,0xC80},{0xC85,0xC8C}, + {0xC8E,0xC90},{0xC92,0xCA8},{0xCAA,0xCB3},{0xCB5,0xCB9},{0xCBD,0xCBD},{0xCDD,0xCDE}, + {0xCE0,0xCE1},{0xCF1,0xCF2},{0xD04,0xD0C},{0xD0E,0xD10},{0xD12,0xD3A},{0xD3D,0xD3D}, + {0xD4E,0xD4E},{0xD54,0xD56},{0xD5F,0xD61},{0xD7A,0xD7F},{0xD85,0xD96},{0xD9A,0xDB1}, + {0xDB3,0xDBB},{0xDBD,0xDBD},{0xDC0,0xDC6},{0xE01,0xE30},{0xE32,0xE33},{0xE40,0xE46}, + {0xE81,0xE82},{0xE84,0xE84},{0xE86,0xE8A},{0xE8C,0xEA3},{0xEA5,0xEA5},{0xEA7,0xEB0}, + {0xEB2,0xEB3},{0xEBD,0xEBD},{0xEC0,0xEC4},{0xEC6,0xEC6},{0xEDC,0xEDF},{0xF00,0xF00}, + {0xF40,0xF47},{0xF49,0xF6C},{0xF88,0xF8C},{0x1000,0x102A},{0x103F,0x103F},{0x1050,0x1055}, + {0x105A,0x105D},{0x1061,0x1061},{0x1065,0x1066},{0x106E,0x1070},{0x1075,0x1081},{0x108E,0x108E}, + {0x10A0,0x10C5},{0x10C7,0x10C7},{0x10CD,0x10CD},{0x10D0,0x10FA},{0x10FC,0x1248},{0x124A,0x124D}, + {0x1250,0x1256},{0x1258,0x1258},{0x125A,0x125D},{0x1260,0x1288},{0x128A,0x128D},{0x1290,0x12B0}, + {0x12B2,0x12B5},{0x12B8,0x12BE},{0x12C0,0x12C0},{0x12C2,0x12C5},{0x12C8,0x12D6},{0x12D8,0x1310}, + {0x1312,0x1315},{0x1318,0x135A},{0x1380,0x138F},{0x13A0,0x13F5},{0x13F8,0x13FD},{0x1401,0x166C}, + {0x166F,0x167F},{0x1681,0x169A},{0x16A0,0x16EA},{0x16F1,0x16F8},{0x1700,0x1711},{0x171F,0x1731}, + {0x1740,0x1751},{0x1760,0x176C},{0x176E,0x1770},{0x1780,0x17B3},{0x17D7,0x17D7},{0x17DC,0x17DC}, + {0x1820,0x1878},{0x1880,0x1884},{0x1887,0x18A8},{0x18AA,0x18AA},{0x18B0,0x18F5},{0x1900,0x191E}, + {0x1950,0x196D},{0x1970,0x1974},{0x1980,0x19AB},{0x19B0,0x19C9},{0x1A00,0x1A16},{0x1A20,0x1A54}, + {0x1AA7,0x1AA7},{0x1B05,0x1B33},{0x1B45,0x1B4C},{0x1B83,0x1BA0},{0x1BAE,0x1BAF},{0x1BBA,0x1BE5}, + {0x1C00,0x1C23},{0x1C4D,0x1C4F},{0x1C5A,0x1C7D},{0x1C80,0x1C88},{0x1C90,0x1CBA},{0x1CBD,0x1CBF}, + {0x1CE9,0x1CEC},{0x1CEE,0x1CF3},{0x1CF5,0x1CF6},{0x1CFA,0x1CFA},{0x1D00,0x1DBF},{0x1E00,0x1F15}, + {0x1F18,0x1F1D},{0x1F20,0x1F45},{0x1F48,0x1F4D},{0x1F50,0x1F57},{0x1F59,0x1F59},{0x1F5B,0x1F5B}, + {0x1F5D,0x1F5D},{0x1F5F,0x1F7D},{0x1F80,0x1FB4},{0x1FB6,0x1FBC},{0x1FBE,0x1FBE},{0x1FC2,0x1FC4}, + {0x1FC6,0x1FCC},{0x1FD0,0x1FD3},{0x1FD6,0x1FDB},{0x1FE0,0x1FEC},{0x1FF2,0x1FF4},{0x1FF6,0x1FFC}, + {0x2071,0x2071},{0x207F,0x207F},{0x2090,0x209C},{0x2102,0x2102},{0x2107,0x2107},{0x210A,0x2113}, + {0x2115,0x2115},{0x2119,0x211D},{0x2124,0x2124},{0x2126,0x2126},{0x2128,0x2128},{0x212A,0x212D}, + {0x212F,0x2139},{0x213C,0x213F},{0x2145,0x2149},{0x214E,0x214E},{0x2183,0x2184},{0x2C00,0x2CE4}, + {0x2CEB,0x2CEE},{0x2CF2,0x2CF3},{0x2D00,0x2D25},{0x2D27,0x2D27},{0x2D2D,0x2D2D},{0x2D30,0x2D67}, + {0x2D6F,0x2D6F},{0x2D80,0x2D96},{0x2DA0,0x2DA6},{0x2DA8,0x2DAE},{0x2DB0,0x2DB6},{0x2DB8,0x2DBE}, + {0x2DC0,0x2DC6},{0x2DC8,0x2DCE},{0x2DD0,0x2DD6},{0x2DD8,0x2DDE},{0x2E2F,0x2E2F},{0x3005,0x3006}, + {0x3031,0x3035},{0x303B,0x303C},{0x3041,0x3096},{0x309D,0x309F},{0x30A1,0x30FA},{0x30FC,0x30FF}, + {0x3105,0x312F},{0x3131,0x318E},{0x31A0,0x31BF},{0x31F0,0x31FF},{0x3400,0x4DBF},{0x4E00,0xA48C}, + {0xA4D0,0xA4FD},{0xA500,0xA60C},{0xA610,0xA61F},{0xA62A,0xA62B},{0xA640,0xA66E},{0xA67F,0xA69D}, + {0xA6A0,0xA6E5},{0xA717,0xA71F},{0xA722,0xA788},{0xA78B,0xA7CA},{0xA7D0,0xA7D1},{0xA7D3,0xA7D3}, + {0xA7D5,0xA7D9},{0xA7F2,0xA801},{0xA803,0xA805},{0xA807,0xA80A},{0xA80C,0xA822},{0xA840,0xA873}, + {0xA882,0xA8B3},{0xA8F2,0xA8F7},{0xA8FB,0xA8FB},{0xA8FD,0xA8FE},{0xA90A,0xA925},{0xA930,0xA946}, + {0xA960,0xA97C},{0xA984,0xA9B2},{0xA9CF,0xA9CF},{0xA9E0,0xA9E4},{0xA9E6,0xA9EF},{0xA9FA,0xA9FE}, + {0xAA00,0xAA28},{0xAA40,0xAA42},{0xAA44,0xAA4B},{0xAA60,0xAA76},{0xAA7A,0xAA7A},{0xAA7E,0xAAAF}, + {0xAAB1,0xAAB1},{0xAAB5,0xAAB6},{0xAAB9,0xAABD},{0xAAC0,0xAAC0},{0xAAC2,0xAAC2},{0xAADB,0xAADD}, + {0xAAE0,0xAAEA},{0xAAF2,0xAAF4},{0xAB01,0xAB06},{0xAB09,0xAB0E},{0xAB11,0xAB16},{0xAB20,0xAB26}, + {0xAB28,0xAB2E},{0xAB30,0xAB5A},{0xAB5C,0xAB69},{0xAB70,0xABE2},{0xAC00,0xD7A3},{0xD7B0,0xD7C6}, + {0xD7CB,0xD7FB},{0xF900,0xFA6D},{0xFA70,0xFAD9},{0xFB00,0xFB06},{0xFB13,0xFB17},{0xFB1D,0xFB1D}, + {0xFB1F,0xFB28},{0xFB2A,0xFB36},{0xFB38,0xFB3C},{0xFB3E,0xFB3E},{0xFB40,0xFB41},{0xFB43,0xFB44}, + {0xFB46,0xFBB1},{0xFBD3,0xFD3D},{0xFD50,0xFD8F},{0xFD92,0xFDC7},{0xFDF0,0xFDFB},{0xFE70,0xFE74}, + {0xFE76,0xFEFC},{0xFF21,0xFF3A},{0xFF41,0xFF5A},{0xFF66,0xFFBE},{0xFFC2,0xFFC7},{0xFFCA,0xFFCF}, + {0xFFD2,0xFFD7},{0xFFDA,0xFFDC},{0x10000,0x1000B},{0x1000D,0x10026},{0x10028,0x1003A},{0x1003C,0x1003D}, + {0x1003F,0x1004D},{0x10050,0x1005D},{0x10080,0x100FA},{0x10280,0x1029C},{0x102A0,0x102D0},{0x10300,0x1031F}, + {0x1032D,0x10340},{0x10342,0x10349},{0x10350,0x10375},{0x10380,0x1039D},{0x103A0,0x103C3},{0x103C8,0x103CF}, + {0x10400,0x1049D},{0x104B0,0x104D3},{0x104D8,0x104FB},{0x10500,0x10527},{0x10530,0x10563},{0x10570,0x1057A}, + {0x1057C,0x1058A},{0x1058C,0x10592},{0x10594,0x10595},{0x10597,0x105A1},{0x105A3,0x105B1},{0x105B3,0x105B9}, + {0x105BB,0x105BC},{0x10600,0x10736},{0x10740,0x10755},{0x10760,0x10767},{0x10780,0x10785},{0x10787,0x107B0}, + {0x107B2,0x107BA},{0x10800,0x10805},{0x10808,0x10808},{0x1080A,0x10835},{0x10837,0x10838},{0x1083C,0x1083C}, + {0x1083F,0x10855},{0x10860,0x10876},{0x10880,0x1089E},{0x108E0,0x108F2},{0x108F4,0x108F5},{0x10900,0x10915}, + {0x10920,0x10939},{0x10980,0x109B7},{0x109BE,0x109BF},{0x10A00,0x10A00},{0x10A10,0x10A13},{0x10A15,0x10A17}, + {0x10A19,0x10A35},{0x10A60,0x10A7C},{0x10A80,0x10A9C},{0x10AC0,0x10AC7},{0x10AC9,0x10AE4},{0x10B00,0x10B35}, + {0x10B40,0x10B55},{0x10B60,0x10B72},{0x10B80,0x10B91},{0x10C00,0x10C48},{0x10C80,0x10CB2},{0x10CC0,0x10CF2}, + {0x10D00,0x10D23},{0x10E80,0x10EA9},{0x10EB0,0x10EB1},{0x10F00,0x10F1C},{0x10F27,0x10F27},{0x10F30,0x10F45}, + {0x10F70,0x10F81},{0x10FB0,0x10FC4},{0x10FE0,0x10FF6},{0x11003,0x11037},{0x11071,0x11072},{0x11075,0x11075}, + {0x11083,0x110AF},{0x110D0,0x110E8},{0x11103,0x11126},{0x11144,0x11144},{0x11147,0x11147},{0x11150,0x11172}, + {0x11176,0x11176},{0x11183,0x111B2},{0x111C1,0x111C4},{0x111DA,0x111DA},{0x111DC,0x111DC},{0x11200,0x11211}, + {0x11213,0x1122B},{0x1123F,0x11240},{0x11280,0x11286},{0x11288,0x11288},{0x1128A,0x1128D},{0x1128F,0x1129D}, + {0x1129F,0x112A8},{0x112B0,0x112DE},{0x11305,0x1130C},{0x1130F,0x11310},{0x11313,0x11328},{0x1132A,0x11330}, + {0x11332,0x11333},{0x11335,0x11339},{0x1133D,0x1133D},{0x11350,0x11350},{0x1135D,0x11361},{0x11400,0x11434}, + {0x11447,0x1144A},{0x1145F,0x11461},{0x11480,0x114AF},{0x114C4,0x114C5},{0x114C7,0x114C7},{0x11580,0x115AE}, + {0x115D8,0x115DB},{0x11600,0x1162F},{0x11644,0x11644},{0x11680,0x116AA},{0x116B8,0x116B8},{0x11700,0x1171A}, + {0x11740,0x11746},{0x11800,0x1182B},{0x118A0,0x118DF},{0x118FF,0x11906},{0x11909,0x11909},{0x1190C,0x11913}, + {0x11915,0x11916},{0x11918,0x1192F},{0x1193F,0x1193F},{0x11941,0x11941},{0x119A0,0x119A7},{0x119AA,0x119D0}, + {0x119E1,0x119E1},{0x119E3,0x119E3},{0x11A00,0x11A00},{0x11A0B,0x11A32},{0x11A3A,0x11A3A},{0x11A50,0x11A50}, + {0x11A5C,0x11A89},{0x11A9D,0x11A9D},{0x11AB0,0x11AF8},{0x11C00,0x11C08},{0x11C0A,0x11C2E},{0x11C40,0x11C40}, + {0x11C72,0x11C8F},{0x11D00,0x11D06},{0x11D08,0x11D09},{0x11D0B,0x11D30},{0x11D46,0x11D46},{0x11D60,0x11D65}, + {0x11D67,0x11D68},{0x11D6A,0x11D89},{0x11D98,0x11D98},{0x11EE0,0x11EF2},{0x11F02,0x11F02},{0x11F04,0x11F10}, + {0x11F12,0x11F33},{0x11FB0,0x11FB0},{0x12000,0x12399},{0x12480,0x12543},{0x12F90,0x12FF0},{0x13000,0x1342F}, + {0x13441,0x13446},{0x14400,0x14646},{0x16800,0x16A38},{0x16A40,0x16A5E},{0x16A70,0x16ABE},{0x16AD0,0x16AED}, + {0x16B00,0x16B2F},{0x16B40,0x16B43},{0x16B63,0x16B77},{0x16B7D,0x16B8F},{0x16E40,0x16E7F},{0x16F00,0x16F4A}, + {0x16F50,0x16F50},{0x16F93,0x16F9F},{0x16FE0,0x16FE1},{0x16FE3,0x16FE3},{0x17000,0x187F7},{0x18800,0x18CD5}, + {0x18D00,0x18D08},{0x1AFF0,0x1AFF3},{0x1AFF5,0x1AFFB},{0x1AFFD,0x1AFFE},{0x1B000,0x1B122},{0x1B132,0x1B132}, + {0x1B150,0x1B152},{0x1B155,0x1B155},{0x1B164,0x1B167},{0x1B170,0x1B2FB},{0x1BC00,0x1BC6A},{0x1BC70,0x1BC7C}, + {0x1BC80,0x1BC88},{0x1BC90,0x1BC99},{0x1D400,0x1D454},{0x1D456,0x1D49C},{0x1D49E,0x1D49F},{0x1D4A2,0x1D4A2}, + {0x1D4A5,0x1D4A6},{0x1D4A9,0x1D4AC},{0x1D4AE,0x1D4B9},{0x1D4BB,0x1D4BB},{0x1D4BD,0x1D4C3},{0x1D4C5,0x1D505}, + {0x1D507,0x1D50A},{0x1D50D,0x1D514},{0x1D516,0x1D51C},{0x1D51E,0x1D539},{0x1D53B,0x1D53E},{0x1D540,0x1D544}, + {0x1D546,0x1D546},{0x1D54A,0x1D550},{0x1D552,0x1D6A5},{0x1D6A8,0x1D6C0},{0x1D6C2,0x1D6DA},{0x1D6DC,0x1D6FA}, + {0x1D6FC,0x1D714},{0x1D716,0x1D734},{0x1D736,0x1D74E},{0x1D750,0x1D76E},{0x1D770,0x1D788},{0x1D78A,0x1D7A8}, + {0x1D7AA,0x1D7C2},{0x1D7C4,0x1D7CB},{0x1DF00,0x1DF1E},{0x1DF25,0x1DF2A},{0x1E030,0x1E06D},{0x1E100,0x1E12C}, + {0x1E137,0x1E13D},{0x1E14E,0x1E14E},{0x1E290,0x1E2AD},{0x1E2C0,0x1E2EB},{0x1E4D0,0x1E4EB},{0x1E7E0,0x1E7E6}, + {0x1E7E8,0x1E7EB},{0x1E7ED,0x1E7EE},{0x1E7F0,0x1E7FE},{0x1E800,0x1E8C4},{0x1E900,0x1E943},{0x1E94B,0x1E94B}, + {0x1EE00,0x1EE03},{0x1EE05,0x1EE1F},{0x1EE21,0x1EE22},{0x1EE24,0x1EE24},{0x1EE27,0x1EE27},{0x1EE29,0x1EE32}, + {0x1EE34,0x1EE37},{0x1EE39,0x1EE39},{0x1EE3B,0x1EE3B},{0x1EE42,0x1EE42},{0x1EE47,0x1EE47},{0x1EE49,0x1EE49}, + {0x1EE4B,0x1EE4B},{0x1EE4D,0x1EE4F},{0x1EE51,0x1EE52},{0x1EE54,0x1EE54},{0x1EE57,0x1EE57},{0x1EE59,0x1EE59}, + {0x1EE5B,0x1EE5B},{0x1EE5D,0x1EE5D},{0x1EE5F,0x1EE5F},{0x1EE61,0x1EE62},{0x1EE64,0x1EE64},{0x1EE67,0x1EE6A}, + {0x1EE6C,0x1EE72},{0x1EE74,0x1EE77},{0x1EE79,0x1EE7C},{0x1EE7E,0x1EE7E},{0x1EE80,0x1EE89},{0x1EE8B,0x1EE9B}, + {0x1EEA1,0x1EEA3},{0x1EEA5,0x1EEA9},{0x1EEAB,0x1EEBB},{0x20000,0x2A6DF},{0x2A700,0x2B739},{0x2B740,0x2B81D}, + {0x2B820,0x2CEA1},{0x2CEB0,0x2EBE0},{0x2F800,0x2FA1D},{0x30000,0x3134A},{0x31350,0x323AF}, +}; +static const int uni_L_n = 659; + +static const uint32_t uni_N[][2] = { + {0x30,0x39},{0xB2,0xB3},{0xB9,0xB9},{0xBC,0xBE},{0x660,0x669},{0x6F0,0x6F9}, + {0x7C0,0x7C9},{0x966,0x96F},{0x9E6,0x9EF},{0x9F4,0x9F9},{0xA66,0xA6F},{0xAE6,0xAEF}, + {0xB66,0xB6F},{0xB72,0xB77},{0xBE6,0xBF2},{0xC66,0xC6F},{0xC78,0xC7E},{0xCE6,0xCEF}, + {0xD58,0xD5E},{0xD66,0xD78},{0xDE6,0xDEF},{0xE50,0xE59},{0xED0,0xED9},{0xF20,0xF33}, + {0x1040,0x1049},{0x1090,0x1099},{0x1369,0x137C},{0x16EE,0x16F0},{0x17E0,0x17E9},{0x17F0,0x17F9}, + {0x1810,0x1819},{0x1946,0x194F},{0x19D0,0x19DA},{0x1A80,0x1A89},{0x1A90,0x1A99},{0x1B50,0x1B59}, + {0x1BB0,0x1BB9},{0x1C40,0x1C49},{0x1C50,0x1C59},{0x2070,0x2070},{0x2074,0x2079},{0x2080,0x2089}, + {0x2150,0x2182},{0x2185,0x2189},{0x2460,0x249B},{0x24EA,0x24FF},{0x2776,0x2793},{0x2CFD,0x2CFD}, + {0x3007,0x3007},{0x3021,0x3029},{0x3038,0x303A},{0x3192,0x3195},{0x3220,0x3229},{0x3248,0x324F}, + {0x3251,0x325F},{0x3280,0x3289},{0x32B1,0x32BF},{0xA620,0xA629},{0xA6E6,0xA6EF},{0xA830,0xA835}, + {0xA8D0,0xA8D9},{0xA900,0xA909},{0xA9D0,0xA9D9},{0xA9F0,0xA9F9},{0xAA50,0xAA59},{0xABF0,0xABF9}, + {0xFF10,0xFF19},{0x10107,0x10133},{0x10140,0x10178},{0x1018A,0x1018B},{0x102E1,0x102FB},{0x10320,0x10323}, + {0x10341,0x10341},{0x1034A,0x1034A},{0x103D1,0x103D5},{0x104A0,0x104A9},{0x10858,0x1085F},{0x10879,0x1087F}, + {0x108A7,0x108AF},{0x108FB,0x108FF},{0x10916,0x1091B},{0x109BC,0x109BD},{0x109C0,0x109CF},{0x109D2,0x109FF}, + {0x10A40,0x10A48},{0x10A7D,0x10A7E},{0x10A9D,0x10A9F},{0x10AEB,0x10AEF},{0x10B58,0x10B5F},{0x10B78,0x10B7F}, + {0x10BA9,0x10BAF},{0x10CFA,0x10CFF},{0x10D30,0x10D39},{0x10E60,0x10E7E},{0x10F1D,0x10F26},{0x10F51,0x10F54}, + {0x10FC5,0x10FCB},{0x11052,0x1106F},{0x110F0,0x110F9},{0x11136,0x1113F},{0x111D0,0x111D9},{0x111E1,0x111F4}, + {0x112F0,0x112F9},{0x11450,0x11459},{0x114D0,0x114D9},{0x11650,0x11659},{0x116C0,0x116C9},{0x11730,0x1173B}, + {0x118E0,0x118F2},{0x11950,0x11959},{0x11C50,0x11C6C},{0x11D50,0x11D59},{0x11DA0,0x11DA9},{0x11F50,0x11F59}, + {0x11FC0,0x11FD4},{0x12400,0x1246E},{0x16A60,0x16A69},{0x16AC0,0x16AC9},{0x16B50,0x16B59},{0x16B5B,0x16B61}, + {0x16E80,0x16E96},{0x1D2C0,0x1D2D3},{0x1D2E0,0x1D2F3},{0x1D360,0x1D378},{0x1D7CE,0x1D7FF},{0x1E140,0x1E149}, + {0x1E2F0,0x1E2F9},{0x1E4F0,0x1E4F9},{0x1E8C7,0x1E8CF},{0x1E950,0x1E959},{0x1EC71,0x1ECAB},{0x1ECAD,0x1ECAF}, + {0x1ECB1,0x1ECB4},{0x1ED01,0x1ED2D},{0x1ED2F,0x1ED3D},{0x1F100,0x1F10C},{0x1FBF0,0x1FBF9}, +}; +static const int uni_N_n = 137; + +static const uint32_t uni_S[][2] = { + {0x9,0xD},{0x20,0x20},{0x85,0x85},{0xA0,0xA0},{0x1680,0x1680},{0x2000,0x200A}, + {0x2028,0x2029},{0x202F,0x202F},{0x205F,0x205F},{0x3000,0x3000}, +}; +static const int uni_S_n = 10; + +static int uni_in(const uint32_t t[][2], int n, uint32_t cp){ + int lo=0, hi=n-1; + while(lo<=hi){ int m=(lo+hi)>>1; + if(cpt[m][1]) lo=m+1; else return 1; } + return 0; +} +static inline int is_L(uint32_t c){ return uni_in(uni_L,uni_L_n,c); } +static inline int is_N(uint32_t c){ return uni_in(uni_N,uni_N_n,c); } +static inline int is_S(uint32_t c){ return uni_in(uni_S,uni_S_n,c); } +#endif diff --git a/engine.py b/engine.py new file mode 100644 index 0000000..31b41fb --- /dev/null +++ b/engine.py @@ -0,0 +1,219 @@ +""" +Motore di inferenza MoE con EXPERT-STREAMING dal disco. + +Idea (quella dell'utente, resa reale): + - la parte DENSA (embedding, attenzione, router, norme, lm_head) sta in RAM; + - gli EXPERT stanno su disco in un file safetensors mappato in memoria (mmap) + e vengono letti SOLO quando un token li attiva; + - una cache LRU tiene in RAM gli expert "caldi" -> meno letture da disco. + +Cosi' un modello che NON entra in RAM gira lo stesso: in RAM ci tieni solo +densa + cache, il resto vive sul disco. Validato qui su OLMoE-1B-7B. + +NB: scritto per OLMoE (Llama-style con QK-norm). I punti specifici del modello +(routing, norme) sono isolati cosi' che lo stesso scheletro valga per GLM/DeepSeek. +""" +import os, json, glob, struct, time, mmap, collections +import torch +import torch.nn.functional as F + +ST_DTYPE = {"BF16": torch.bfloat16, "F16": torch.float16, "F32": torch.float32} + + +class Shards: + """Indicizza i tensori in piu' file safetensors e li legge via mmap on-demand.""" + def __init__(self, snap_dir): + self.index = {} # name -> (shard_path, abs_offset, nbytes, torch_dtype, shape) + self.mm = {} # shard_path -> mmap + for shard in sorted(glob.glob(os.path.join(snap_dir, "model-*.safetensors"))): + with open(shard, "rb") as f: + hlen = struct.unpack("= 16: + return w + qmax = (1 << (bits - 1)) - 1 # int8->127, int4->7, int3->3, int2->1 + wf = w.float() + scale = wf.abs().amax(dim=1, keepdim=True) / qmax + scale = scale.clamp_min(1e-8) + wq = torch.round(wf / scale).clamp(-qmax - 1, qmax) + return (wq * scale).to(torch.bfloat16) + + +class ExpertCache: + """Cache LRU degli expert. capacity = quanti expert teniamo residenti PER LAYER.""" + def __init__(self, shards, n_layers, capacity, quant_bits=16): + self.shards = shards + self.cap = capacity + self.quant_bits = quant_bits + self.caches = [collections.OrderedDict() for _ in range(n_layers)] + self.hits = 0 + self.miss = 0 + + def get(self, layer, eid): + """Ritorna (gate_w, up_w, down_w) dell'expert, da cache o da disco.""" + c = self.caches[layer] + if eid in c: + self.hits += 1 + c.move_to_end(eid) + return c[eid] + self.miss += 1 + p = f"model.layers.{layer}.mlp.experts.{eid}." + # tengo gli expert in bf16 (niente .float(): -24% tempo, -50% RAM, piu' fedele al riferimento) + b = self.quant_bits + w = (quant_dequant(self.shards.read(p + "gate_proj.weight"), b), + quant_dequant(self.shards.read(p + "up_proj.weight"), b), + quant_dequant(self.shards.read(p + "down_proj.weight"), b)) + c[eid] = w + if len(c) > self.cap: + c.popitem(last=False) + return w + + def hitrate(self): + t = self.hits + self.miss + return self.hits / t if t else 0.0 + + +def rmsnorm(x, w, eps=1e-5): + x = x.float() + x = x * torch.rsqrt(x.pow(2).mean(-1, keepdim=True) + eps) + return x * w.float() + + +def rotate_half(x): + h = x.shape[-1] // 2 + return torch.cat((-x[..., h:], x[..., :h]), dim=-1) + + +class OlmoeStreaming: + def __init__(self, snap_dir, expert_cap=16, quant_bits=16): + self.cfg = json.load(open(os.path.join(snap_dir, "config.json"))) + self.shards = Shards(snap_dir) + c = self.cfg + self.L = c["num_hidden_layers"] + self.H = c["num_attention_heads"] + self.hd = c["hidden_size"] // self.H + self.topk = c["num_experts_per_tok"] + self.eps = c.get("rms_norm_eps", 1e-5) + self.norm_topk = c.get("norm_topk_prob", False) + theta = c.get("rope_theta", 10000.0) + self.inv_freq = 1.0 / (theta ** (torch.arange(0, self.hd, 2).float() / self.hd)) + self.cache = ExpertCache(self.shards, self.L, expert_cap, quant_bits) + + # --- parte DENSA: residente in RAM (float32) --- + t = time.time() + self.embed = self.shards.read("model.embed_tokens.weight").float() + self.lm_head = self.shards.read("lm_head.weight").float() + self.final_norm = self.shards.read("model.norm.weight").float() + self.layers = [] + for i in range(self.L): + p = f"model.layers.{i}." + self.layers.append({ + "in_ln": self.shards.read(p + "input_layernorm.weight").float(), + "post_ln":self.shards.read(p + "post_attention_layernorm.weight").float(), + "q": self.shards.read(p + "self_attn.q_proj.weight").float(), + "k": self.shards.read(p + "self_attn.k_proj.weight").float(), + "v": self.shards.read(p + "self_attn.v_proj.weight").float(), + "o": self.shards.read(p + "self_attn.o_proj.weight").float(), + "qn": self.shards.read(p + "self_attn.q_norm.weight").float(), + "kn": self.shards.read(p + "self_attn.k_norm.weight").float(), + "gate": self.shards.read(p + "mlp.gate.weight").float(), + }) + self.dense_load_s = time.time() - t + + def _rope(self, x, pos): + # x: (heads, seq, hd) ; pos: (seq,) + freqs = torch.outer(pos.float(), self.inv_freq) # (seq, hd/2) + emb = torch.cat((freqs, freqs), dim=-1) # (seq, hd) + cos, sin = emb.cos(), emb.sin() + return x * cos + rotate_half(x) * sin + + def _attn(self, lw, x, pos, layer, kv): + """Attenzione con KV-cache. x = SOLO i token nuovi (S in prefill, 1 in decode). + pos = posizioni assolute dei token nuovi. kv = lista per-layer dei (k,v) passati.""" + S = x.shape[0] + q = rmsnorm(x @ lw["q"].T, lw["qn"], self.eps).view(S, self.H, self.hd).transpose(0, 1) + k = rmsnorm(x @ lw["k"].T, lw["kn"], self.eps).view(S, self.H, self.hd).transpose(0, 1) + v = (x @ lw["v"].T).view(S, self.H, self.hd).transpose(0, 1) + q = self._rope(q, pos); k = self._rope(k, pos) + if kv is not None and kv[layer] is not None: # concateno il passato + pk, pv = kv[layer] + k = torch.cat([pk, k], dim=1); v = torch.cat([pv, v], dim=1) + if kv is not None: + kv[layer] = (k, v) + Tk = k.shape[1] # lunghezza totale (passato+nuovi) + scores = (q @ k.transpose(-1, -2)) / (self.hd ** 0.5) # (H,S,Tk) + # mask causale: query a posizione assoluta pos[i] vede key j<=pos[i] + kpos = torch.arange(Tk) + mask = torch.where(kpos[None, :] > pos[:, None], float("-inf"), 0.0) # -inf dove vietato + a = F.softmax(scores + mask, dim=-1) + out = (a @ v).transpose(0, 1).reshape(S, self.H * self.hd) + return out @ lw["o"].T + + def _moe(self, layer, lw, x): + S = x.shape[0] + logits = x @ lw["gate"].T # (S,64) + probs = F.softmax(logits.float(), dim=-1) + w, idx = torch.topk(probs, self.topk, dim=-1) # (S,topk) + if self.norm_topk: + w = w / w.sum(-1, keepdim=True) + out = torch.zeros_like(x) + # raggruppo per expert: per ogni expert davvero usato, processo i suoi token + for eid in torch.unique(idx).tolist(): + sel = (idx == eid) # (S,topk) bool + rows = sel.any(dim=-1).nonzero(as_tuple=True)[0] + if rows.numel() == 0: + continue + gw, uw, dw = self.cache.get(layer, eid) # <-- streaming dal disco (bf16) + xe = x[rows].to(torch.bfloat16) # calcolo expert in bf16 + h = (F.silu(xe @ gw.T) * (xe @ uw.T)) @ dw.T + weight = (w[rows] * sel[rows].float()).sum(-1, keepdim=True) + out[rows] += weight * h.float() + return out + + @torch.no_grad() + def _step(self, ids_new, pos, kv): + """Un passo del modello sui token nuovi. Ritorna logit dell'ultimo token.""" + x = self.embed[torch.tensor(ids_new)] # (S,hidden) + for i, lw in enumerate(self.layers): + x = x + self._attn(lw, rmsnorm(x, lw["in_ln"], self.eps), pos, i, kv) + x = x + self._moe(i, lw, rmsnorm(x, lw["post_ln"], self.eps)) + x = rmsnorm(x, self.final_norm, self.eps) + return (x @ self.lm_head.T)[-1] + + @torch.no_grad() + def generate(self, token_ids, n_new, greedy=True): + kv = [None] * self.L + ids = list(token_ids) + # PREFILL: tutti i token del prompt in un colpo, riempie la kv-cache + logit = self._step(ids, torch.arange(len(ids)), kv) + for s in range(n_new): + nxt = int(torch.argmax(logit)) if greedy else int(torch.multinomial(F.softmax(logit, -1), 1)) + ids.append(nxt) + if s == n_new - 1: + break + # DECODE: un solo token nuovo, usa la kv-cache (qui la cache expert torna a funzionare) + logit = self._step([nxt], torch.tensor([len(ids) - 1]), kv) + return ids diff --git a/profile_run.py b/profile_run.py new file mode 100644 index 0000000..b821850 --- /dev/null +++ b/profile_run.py @@ -0,0 +1,18 @@ +"""Profila dove va il tempo: lettura expert dal disco vs attenzione vs moe vs matmul.""" +import cProfile, pstats, io, glob, json +from engine import OlmoeStreaming + +snap = glob.glob("/home/vincenzo/.cache/huggingface/hub/models--allenai--OLMoE-1B-7B-0924/snapshots/*")[0] +ref = json.load(open("ref.json")) +m = OlmoeStreaming(snap, expert_cap=16) + +pr = cProfile.Profile() +pr.enable() +m.generate(ref["prompt_ids"], 8, greedy=True) # 8 token bastano per il profilo +pr.disable() + +s = io.StringIO() +ps = pstats.Stats(pr, stream=s).sort_stats("tottime") +ps.print_stats(15) +print(s.getvalue()) +print(f"Hit-rate: {m.cache.hitrate()*100:.1f}% hit={m.cache.hits} miss={m.cache.miss}") diff --git a/quant_test.py b/quant_test.py new file mode 100644 index 0000000..2a6f7f8 --- /dev/null +++ b/quant_test.py @@ -0,0 +1,22 @@ +"""La domanda che conta: a quanti bit l'output degli expert REGGE ancora? +Quantizzo solo gli expert (la parte densa resta bf16) e confronto col riferimento.""" +import json, glob +from engine import OlmoeStreaming + +snap = glob.glob("/home/vincenzo/.cache/huggingface/hub/models--allenai--OLMoE-1B-7B-0924/snapshots/*")[0] +ref = json.load(open("ref.json")) +exp = ref["full_ids"][len(ref["prompt_ids"]):] +n_new = len(exp) + +print(f"{'bit':>4} {'MB/expert':>10} {'match':>7} testo") +for bits in (16, 8, 4, 3, 2): + m = OlmoeStreaming(snap, expert_cap=64, quant_bits=bits) # cap64: isola l'effetto quant + out = m.generate(ref["prompt_ids"], n_new, greedy=True) + gen = out[len(ref["prompt_ids"]):] + match = sum(a == b for a, b in zip(gen, exp)) + mb = 6.29 * bits / 8 / 1.0 # ~6.29M param/expert * bit / 8 -> MB + # decode testo per vedere se e' ancora sensato + from transformers import AutoTokenizer + tok = AutoTokenizer.from_pretrained("allenai/OLMoE-1B-7B-0924") + txt = tok.decode(gen) + print(f"{bits:>4} {mb:>9.1f}M {match:>4}/{n_new:<2} {txt!r}") diff --git a/ref.json b/ref.json new file mode 100644 index 0000000..7c236ab --- /dev/null +++ b/ref.json @@ -0,0 +1 @@ +{"prompt": "The capital of France is", "prompt_ids": [510, 5347, 273, 6181, 310], "full_ids": [510, 5347, 273, 6181, 310, 7785, 15, 187, 187, 510, 5347, 273, 253, 1986, 2077, 310, 5041], "text": "The capital of France is Paris.\n\nThe capital of the United States is Washington"} \ No newline at end of file diff --git a/run.py b/run.py new file mode 100644 index 0000000..c4acda0 --- /dev/null +++ b/run.py @@ -0,0 +1,31 @@ +"""Lancia il motore streaming, confronta con il riferimento, misura RAM/hit-rate/velocita'.""" +import json, time, glob, sys, resource +from engine import OlmoeStreaming + +CAP = int(sys.argv[1]) if len(sys.argv) > 1 else 16 # expert residenti per layer (su 64) +snap = glob.glob("/home/vincenzo/.cache/huggingface/hub/models--allenai--OLMoE-1B-7B-0924/snapshots/*")[0] +ref = json.load(open("run_ref.json" if False else "ref.json")) + +def rss_gb(): + return resource.getrusage(resource.RUSAGE_SELF).ru_maxrss / (1024**2) # KB->GB su Linux + +print(f"== Motore streaming, cache = {CAP} expert/layer su 64 ==") +t = time.time() +m = OlmoeStreaming(snap, expert_cap=CAP) +print(f"densa caricata in {m.dense_load_s:.1f}s | RSS dopo load densa: {rss_gb():.2f} GB") + +n_new = len(ref["full_ids"]) - len(ref["prompt_ids"]) +t = time.time() +out = m.generate(ref["prompt_ids"], n_new, greedy=True) +dt = time.time() - t + +# confronto +gen = out[len(ref["prompt_ids"]):] +exp = ref["full_ids"][len(ref["prompt_ids"]):] +match = sum(a == b for a, b in zip(gen, exp)) +print(f"\nRiferimento (transformers): {exp}") +print(f"Motore streaming : {gen}") +print(f"Token coincidenti: {match}/{len(exp)}") +print(f"\nRSS PICCO: {rss_gb():.2f} GB (modello completo in bf16 = ~13 GB)") +print(f"Hit-rate cache expert: {m.cache.hitrate()*100:.1f}% (hit={m.cache.hits} miss={m.cache.miss})") +print(f"Velocita': {n_new/dt:.2f} tok/s ({dt:.1f}s per {n_new} token, no kv-cache)") diff --git a/s0_costmodel.py b/s0_costmodel.py new file mode 100644 index 0000000..160ef76 --- /dev/null +++ b/s0_costmodel.py @@ -0,0 +1,97 @@ +""" +Stadio 0 - Cost model + benchmark del disco per lo streaming degli expert MoE. + +Domanda: e' FISICAMENTE possibile fare streaming degli expert da disco +e generare a velocita' usabile, su QUESTA macchina? + +Due numeri che servono: + 1. Banda effettiva del disco in lettura RANDOM, a blocchi grossi quanto un expert. + 2. Quanti byte/token dobbiamo leggere -> da cui il tetto di token/sec. + +Nessun modello richiesto. Gira in secondi. +""" +import os, sys, time, mmap, random, argparse + +MB = 1024 * 1024 +GB = 1024 * MB + + +def bench_disk(path_dir, expert_mb=12.0, total_mb=2048, n_reads=200): + """Crea un file, poi misura lettura sequenziale e random a chunk = un expert.""" + os.makedirs(path_dir, exist_ok=True) + fpath = os.path.join(path_dir, "_bench.bin") + chunk = int(expert_mb * MB) + total = int(total_mb * MB) + total = (total // chunk) * chunk + n_chunks = total // chunk + + # scrittura + t = time.time() + with open(fpath, "wb") as f: + buf = os.urandom(chunk) + for _ in range(n_chunks): + f.write(buf) + f.flush(); os.fsync(f.fileno()) + write_bw = total / (time.time() - t) / GB + + # prova a buttare via la page cache (best effort, serve permessi su Linux nativo) + try: + os.system("sync") + with open("/proc/sys/vm/drop_caches", "w") as c: + c.write("3") + except Exception: + pass # su /mnt/c o senza root non si puo': il numero sara' ottimistico + + f = open(fpath, "rb") + mm = mmap.mmap(f.fileno(), 0, prot=mmap.PROT_READ) + + # random reads a chunk di un expert + idxs = [random.randrange(n_chunks) for _ in range(n_reads)] + s = 0 + t = time.time() + for i in idxs: + off = i * chunk + s += mm[off] # tocca la prima pagina + s += mm[off + chunk - 1] # e l'ultima -> forza il caricamento del range + _ = bytes(mm[off:off + chunk]) # legge davvero l'intero expert + rand_bw = (n_reads * chunk) / (time.time() - t) / GB + + mm.close(); f.close() + os.remove(fpath) + return write_bw, rand_bw + + +def cost_model(name, n_layers, n_active, expert_mb, disk_bw_gbs, ram_resident_gb): + """Stampa il tetto di token/sec in funzione dell'hit-rate della cache.""" + bytes_cold = n_layers * n_active * expert_mb / 1024 # GB letti per token se 0 cache + print(f"\n--- {name} ---") + print(f" layer={n_layers} expert_attivi/layer={n_active} expert={expert_mb:.1f} MB") + print(f" parte densa residente in RAM stimata: ~{ram_resident_gb:.1f} GB") + print(f" byte da streammare per token (cache fredda): {bytes_cold*1024:.0f} MB") + print(f" tetto token/sec @ banda {disk_bw_gbs:.2f} GB/s, al variare dell'hit-rate cache:") + for hit in (0.0, 0.5, 0.8, 0.9, 0.95, 0.99): + eff = bytes_cold * (1 - hit) + tps = disk_bw_gbs / eff if eff > 0 else float("inf") + print(f" hit {hit*100:5.1f}% -> {tps:6.2f} tok/s") + + +if __name__ == "__main__": + ap = argparse.ArgumentParser() + ap.add_argument("--dir", default=".", help="cartella su cui benchmarkare il disco") + ap.add_argument("--expert-mb", type=float, default=12.0) + args = ap.parse_args() + + print(f"Benchmark disco su: {os.path.abspath(args.dir)} (chunk={args.expert_mb} MB)") + wbw, rbw = bench_disk(args.dir, expert_mb=args.expert_mb) + print(f" scrittura seq : {wbw:.2f} GB/s") + print(f" lettura random: {rbw:.2f} GB/s <-- numero che conta per lo streaming") + + # Scenari. expert_mb a Q4 ~ (hidden*inter*3)*0.5B. + # OLMoE 1B-7B: 16 layer, 8 attivi, hidden 2048 inter 1024 -> ~3 MB Q4 + cost_model("OLMoE 1B-7B (piccolo, lo useremo allo Stadio 1)", + n_layers=16, n_active=8, expert_mb=3.0, + disk_bw_gbs=rbw, ram_resident_gb=1.0) + # DeepSeek-V3/V4 class: ~60 layer MoE, 8 attivi, expert ~6 MB Q2 + cost_model("DeepSeek/GLM class @ Q2 (il sogno finale)", + n_layers=60, n_active=8, expert_mb=6.0, + disk_bw_gbs=rbw, ram_resident_gb=10.0) diff --git a/s1_trace_hitrate.py b/s1_trace_hitrate.py new file mode 100644 index 0000000..11c884b --- /dev/null +++ b/s1_trace_hitrate.py @@ -0,0 +1,126 @@ +""" +Pilastro 2 - La domanda che decide tutto: +quanto e' alto l'hit-rate di una cache di expert su un router MoE VERO? + +Facciamo girare OLMoE-1B-7B su testo reale, registriamo per ogni (layer, posizione) +quali 8 expert su 64 vengono attivati, e poi SIMULIAMO diverse politiche di cache +al variare della capacita' K (quanti expert/layer teniamo residenti in RAM). + +Output: hit-rate per policy e per K -> mappato su token/sec col cost model. +""" +import sys, time, collections +import torch + +MODEL = "allenai/OLMoE-1B-7B-0924" +TOPK = 8 +N_EXPERTS = 64 +N_LAYERS = 16 +EXPERT_MB = 12.6 # bf16 + +# Banda misurata allo Stadio 0 (random read, ext4). Aggiorna se rifai il bench. +DISK_BW_GBS = 7.33 + +PROMPTS = [ + "The history of the Roman Empire spans over a thousand years, from the founding " + "of the city to the fall of Constantinople. Its legacy in law, language and " + "engineering still shapes the modern world.", + "def quicksort(arr):\n if len(arr) <= 1:\n return arr\n pivot = arr[len(arr)//2]\n" + " left = [x for x in arr if x < pivot]\n middle = [x for x in arr if x == pivot]\n" + " right = [x for x in arr if x > pivot]\n return quicksort(left) + middle + quicksort(right)", + "In quantum mechanics, the wave function encodes the probability amplitude of a " + "particle's state. Measurement collapses this superposition into a definite outcome, " + "a process that remains philosophically contested.", + "La politica monetaria della banca centrale influenza i tassi di interesse, " + "l'inflazione e l'occupazione. Alzare i tassi raffredda la domanda ma rischia " + "di rallentare la crescita economica.", +] + + +def collect_trace(): + from transformers import AutoModelForCausalLM, AutoTokenizer + print("Carico il modello (bf16, CPU)...", flush=True) + t = time.time() + tok = AutoTokenizer.from_pretrained(MODEL) + model = AutoModelForCausalLM.from_pretrained( + MODEL, torch_dtype=torch.bfloat16, low_cpu_mem_usage=True) + model.eval() + print(f" caricato in {time.time()-t:.0f}s", flush=True) + + # trace[layer] = lista (in ordine di token) di tuple di 8 expert id + trace = [[] for _ in range(N_LAYERS)] + for pi, p in enumerate(PROMPTS): + ids = tok(p, return_tensors="pt").input_ids + with torch.no_grad(): + out = model(ids, output_router_logits=True, use_cache=False) + # out.router_logits: tupla di N_LAYERS tensori (n_token, n_experts) + for li, rl in enumerate(out.router_logits): + topk = rl.topk(TOPK, dim=-1).indices # (n_token, 8) + for row in topk.tolist(): + trace[li].append(tuple(sorted(row))) + print(f" prompt {pi}: {ids.shape[1]} token", flush=True) + return trace + + +def simulate(trace, K, policy="lru"): + """Cache per-layer di capacita' K. Ritorna hit-rate globale sugli accessi a expert.""" + hits = total = 0 + for li in range(N_LAYERS): + if policy == "lru": + cache = collections.OrderedDict() + elif policy == "lfu": + cache = {} # eid -> freq + freq = collections.Counter() + for experts in trace[li]: + for e in experts: + total += 1 + if policy == "lru": + if e in cache: + hits += 1 + cache.move_to_end(e) + else: + cache[e] = 1 + if len(cache) > K: + cache.popitem(last=False) + elif policy == "lfu": + freq[e] += 1 + if e in cache: + hits += 1 + else: + if len(cache) >= K: + victim = min(cache, key=lambda x: freq[x]) + del cache[victim] + cache[e] = 1 + return hits / total if total else 0.0 + + +def consecutive_reuse(trace): + """Frazione di expert al token t gia' attivi al token t-1 (stesso layer).""" + same = tot = 0 + for li in range(N_LAYERS): + seq = trace[li] + for t in range(1, len(seq)): + prev = set(seq[t-1]); cur = set(seq[t]) + same += len(prev & cur); tot += len(cur) + return same / tot if tot else 0.0 + + +def tok_per_sec(hitrate): + bytes_cold_gb = N_LAYERS * TOPK * EXPERT_MB / 1024 + eff = bytes_cold_gb * (1 - hitrate) + return DISK_BW_GBS / eff if eff > 0 else float("inf") + + +if __name__ == "__main__": + trace = collect_trace() + ntok = sum(len(trace[0]) for _ in [0]) + print(f"\nToken totali tracciati: {len(trace[0])} x {N_LAYERS} layer") + print(f"Riuso consecutivo (expert in comune t vs t-1): {consecutive_reuse(trace)*100:.1f}%") + + print("\nHit-rate cache per-layer al variare di K (expert residenti su 64):") + print(f"{'K':>4} {'RAM/GB':>7} {'LRU':>8} {'LFU':>8} {'tok/s@LRU':>10}") + for K in (8, 12, 16, 24, 32, 48): + ram = K * N_LAYERS * EXPERT_MB / 1024 + hl = simulate(trace, K, "lru") + hf = simulate(trace, K, "lfu") + print(f"{K:>4} {ram:>6.1f}G {hl*100:>7.1f}% {hf*100:>7.1f}% {tok_per_sec(hl):>9.1f}") + print("\nNota: K=8 e' il minimo teorico (8 attivi/token). K=64 = tutto in RAM (no streaming).") diff --git a/s2_research.py b/s2_research.py new file mode 100644 index 0000000..00bb90a --- /dev/null +++ b/s2_research.py @@ -0,0 +1,110 @@ +""" +RICERCA - Il cardine del metodo: lo SKEW degli expert e' sfruttabile? + +Se pochi expert "caldi" coprono gran parte delle attivazioni, allora la strategia +giusta per un modello che NON entra in RAM e': + - PIN dei caldi (residenti per sempre in RAM, profilati offline) + - STREAM dei freddi dal disco +invece di una LRU dinamica (che su RAM piccola va in pressione, l'abbiamo visto). + +Test onesto: determino il "set caldo" dalla PRIMA meta' dei token, e misuro la +copertura sulla SECONDA meta' (mai vista). Confronto PIN-caldi statico vs LRU a parita' di K. +""" +import json, glob, collections, time +import torch + +MODEL = "allenai/OLMoE-1B-7B-0924" +N_EXP, TOPK, N_LAYERS = 64, 8, 16 + +# testo piu' lungo e vario per statistiche decenti +PROMPTS = [ + "The Roman Empire was one of the largest empires in history. At its height under " + "Trajan, it covered five million square kilometres and held seventy million people, " + "about a fifth of the world's population at the time. The empire's longevity and vast " + "extent ensured a lasting influence on language, religion, architecture, philosophy, law " + "and forms of government across the territory it once governed. ", + "Photosynthesis is a biological process used by plants, algae and some bacteria to " + "convert light energy into chemical energy stored in glucose. It occurs in the chloroplasts, " + "specifically using the green pigment chlorophyll. The process consumes carbon dioxide and " + "water and releases oxygen as a by-product, sustaining most life on Earth. ", + "def fibonacci(n):\n a, b = 0, 1\n result = []\n for _ in range(n):\n " + "result.append(a)\n a, b = b, a + b\n return result\n\n" + "class Stack:\n def __init__(self):\n self.items = []\n def push(self, x):\n" + " self.items.append(x)\n def pop(self):\n return self.items.pop()\n", + "L'economia mondiale nel ventunesimo secolo e' caratterizzata da una crescente " + "globalizzazione, dall'integrazione dei mercati finanziari e dalla rapida diffusione " + "delle tecnologie digitali. Le banche centrali giocano un ruolo cruciale nel mantenere " + "la stabilita' dei prezzi attraverso la politica monetaria. ", +] + + +def collect(): + from transformers import AutoModelForCausalLM, AutoTokenizer + print("carico modello...", flush=True) + tok = AutoTokenizer.from_pretrained(MODEL) + model = AutoModelForCausalLM.from_pretrained(MODEL, torch_dtype=torch.bfloat16, + low_cpu_mem_usage=True).eval() + trace = [[] for _ in range(N_LAYERS)] + for p in PROMPTS: + ids = tok(p, return_tensors="pt").input_ids + with torch.no_grad(): + out = model(ids, output_router_logits=True, use_cache=False) + for li, rl in enumerate(out.router_logits): + for row in rl.topk(TOPK, -1).indices.tolist(): + trace[li].append(tuple(row)) + print(f" +{ids.shape[1]} token", flush=True) + return trace + + +def lru_hit(seq, K): + c = collections.OrderedDict(); hit = tot = 0 + for experts in seq: + for e in experts: + tot += 1 + if e in c: hit += 1; c.move_to_end(e) + else: + c[e] = 1 + if len(c) > K: c.popitem(last=False) + return hit / tot + + +def static_hot_hit(train, test, K): + """Set caldo = K piu' frequenti nel train; copertura misurata sul test.""" + freq = collections.Counter(e for experts in train for e in experts) + hot = set(e for e, _ in freq.most_common(K)) + hit = tot = 0 + for experts in test: + for e in experts: + tot += 1 + if e in hot: hit += 1 + return hit / tot + + +if __name__ == "__main__": + trace = collect() + ntok = len(trace[0]) + print(f"\nToken totali: {ntok} x {N_LAYERS} layer = {ntok*N_LAYERS*TOPK} accessi expert\n") + + # skew: distribuzione di frequenza (media sui layer), e curva di copertura top-K + print("COPERTURA del set caldo (statico, profilato su prima meta', testato su seconda):") + print(f"{'K':>4} {'RAM':>7} {'pin-caldo':>10} {'LRU':>8} (uniforme=K/64)") + for K in (8, 12, 16, 24, 32): + cov_static, cov_lru = [], [] + for li in range(N_LAYERS): + seq = trace[li]; h = len(seq) // 2 + cov_static.append(static_hot_hit(seq[:h], seq[h:], K)) + cov_lru.append(lru_hit(seq, K)) + cs = sum(cov_static)/N_LAYERS; cl = sum(cov_lru)/N_LAYERS + ram = K * N_LAYERS * 12.6 / 1024 + print(f"{K:>4} {ram:>5.1f}GB {cs*100:>9.1f}% {cl*100:>7.1f}% {K/64*100:>4.0f}%") + + # quanto e' skewata la distribuzione? entropia normalizzata e top-8 share + import math + shares = [] + for li in range(N_LAYERS): + freq = collections.Counter(e for ex in trace[li] for e in ex) + tot = sum(freq.values()) + top8 = sum(c for _, c in freq.most_common(8)) / tot + shares.append(top8) + print(f"\nSkew: gli 8 expert piu' caldi (su 64) coprono in media " + f"{sum(shares)/len(shares)*100:.1f}% delle attivazioni (uniforme sarebbe 12.5%).") diff --git a/sweep.py b/sweep.py new file mode 100644 index 0000000..d088b7a --- /dev/null +++ b/sweep.py @@ -0,0 +1,20 @@ +"""Sweep della cache: per ogni capacita' misura correttezza, hit-rate, RAM cache, velocita'.""" +import json, time, glob +from engine import OlmoeStreaming + +snap = glob.glob("/home/vincenzo/.cache/huggingface/hub/models--allenai--OLMoE-1B-7B-0924/snapshots/*")[0] +ref = json.load(open("ref.json")) +exp = ref["full_ids"][len(ref["prompt_ids"]):] +n_new = len(exp) +EXPERT_MB_BF16 = 12.6 + +print(f"{'cap':>4} {'RAMcache':>9} {'match':>6} {'hit%':>6} {'tok/s':>7} {'sec':>6}") +for cap in (16, 32, 48, 64): + m = OlmoeStreaming(snap, expert_cap=cap) + t = time.time() + out = m.generate(ref["prompt_ids"], n_new, greedy=True) + dt = time.time() - t + gen = out[len(ref["prompt_ids"]):] + match = sum(a == b for a, b in zip(gen, exp)) + ram = cap * m.L * EXPERT_MB_BF16 / 1024 + print(f"{cap:>4} {ram:>7.1f}GB {match:>3}/{n_new:<2} {m.cache.hitrate()*100:>5.1f}% {n_new/dt:>7.2f} {dt:>6.1f}") diff --git a/validate_ref.py b/validate_ref.py new file mode 100644 index 0000000..9c26bb9 --- /dev/null +++ b/validate_ref.py @@ -0,0 +1,17 @@ +"""Genera il riferimento con transformers (greedy) e lo salva. Va lanciato da solo (usa ~13GB).""" +import json, torch +from transformers import AutoModelForCausalLM, AutoTokenizer + +PROMPT = "The capital of France is" +N = 12 +tok = AutoTokenizer.from_pretrained("allenai/OLMoE-1B-7B-0924") +ids = tok(PROMPT, return_tensors="pt").input_ids +model = AutoModelForCausalLM.from_pretrained("allenai/OLMoE-1B-7B-0924", + torch_dtype=torch.bfloat16, low_cpu_mem_usage=True).eval() +with torch.no_grad(): + out = model.generate(ids, max_new_tokens=N, do_sample=False) +full = out[0].tolist() +json.dump({"prompt": PROMPT, "prompt_ids": ids[0].tolist(), "full_ids": full, + "text": tok.decode(full)}, open("ref.json", "w")) +print("RIFERIMENTO salvato:", repr(tok.decode(full))) +print("full_ids:", full)