colibrì: pure-C GLM-5.2 (744B MoE) engine with disk-streamed experts
Engine (c/glm.c): MLA attention with compressed KV, sigmoid noaux_tc router, int8/int4/int2 quant kernels (AVX2), per-layer LRU expert cache + pinned hot-store, batch-union MoE, native MTP speculative decoding (lossless), multi-stop + official chat template, RAM auto-budget from MemAvailable. Tokenizer: byte-level BPE in C. Tooling: coli CLI, disk-safe FP8→int4 converter, tiny-random oracle validation (TF 32/32, greedy 20/20). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
This commit is contained in:
JustVugg
+22
@@ -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
|
||||
@@ -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=<dir> 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+1<len(SPRITE) else "."*len(SPRITE[y])
|
||||
row=""
|
||||
for x in range(len(top)):
|
||||
ct, cb = PAL.get(top[x]), PAL.get(bot[x])
|
||||
if ct is None and cb is None: row+= "\033[0m "
|
||||
elif ct is not None and cb is None: row+= f"\033[38;5;{ct}m\033[49m▀"
|
||||
elif ct is None and cb is not None: row+= f"\033[38;5;{cb}m\033[49m▄"
|
||||
else: row+= f"\033[38;5;{ct}m\033[48;5;{cb}m▀"
|
||||
out.append(row+"\033[0m")
|
||||
return out
|
||||
|
||||
def banner(sub=""):
|
||||
sp=sprite_lines()
|
||||
txt=[
|
||||
f"{C.teal}{C.b}colibrì{C.r} {C.dim}v1.0{C.r}",
|
||||
f"{C.dim}piccolo motore, modello immenso{C.r}",
|
||||
f"{C.gray}GLM-5.2 · 744B MoE · int4 · streaming CPU{C.r}",
|
||||
f"{C.dgray}{sub}{C.r}" if sub else "",
|
||||
"",
|
||||
]
|
||||
print()
|
||||
for i,s in enumerate(sp):
|
||||
t = txt[i] if i<len(txt) else ""
|
||||
print(f" {s} {t}")
|
||||
print(f" {C.dgray}{'─'*58}{C.r}")
|
||||
|
||||
def hline(w): return f"{C.dgray}{'─'*w}{C.r}"
|
||||
|
||||
# ---------- util ----------
|
||||
def term_w(): return min(shutil.get_terminal_size((80,20)).columns, 100)
|
||||
|
||||
def need_model(model):
|
||||
if not os.path.isdir(model):
|
||||
sys.exit(f"{C.yel}modello non trovato:{C.r} {model}\n imposta COLI_MODEL o usa --model")
|
||||
if not os.path.exists(os.path.join(model,"tokenizer.json")):
|
||||
sys.exit(f"{C.yel}manca tokenizer.json in {model}{C.r}")
|
||||
if not os.path.exists(GLM):
|
||||
sys.exit(f"{C.yel}motore non compilato.{C.r} Esegui: coli build")
|
||||
|
||||
def env_for(a):
|
||||
e = dict(os.environ, SNAP=a.model)
|
||||
if a.ram: e["RAM_GB"]=str(a.ram)
|
||||
if a.ngen: e["NGEN"]=str(a.ngen)
|
||||
if a.topp: e["TOPP"]=str(a.topp)
|
||||
if a.topk: e["TOPK"]=str(a.topk)
|
||||
return e
|
||||
|
||||
class Spinner:
|
||||
FRAMES=["⠋","⠙","⠹","⠸","⠼","⠴","⠦","⠧","⠇","⠏"]
|
||||
def __init__(self,label): self.label=label; self.stop_evt=threading.Event(); self.t0=time.time(); self.th=None
|
||||
def start(self):
|
||||
if not TTY: return
|
||||
def run():
|
||||
i=0
|
||||
while not self.stop_evt.is_set():
|
||||
el=time.time()-self.t0
|
||||
sys.stdout.write(f"\r {C.teal}{self.FRAMES[i%10]}{C.r} {C.dim}{self.label} {el:.0f}s{C.r}\033[K")
|
||||
sys.stdout.flush(); i+=1; time.sleep(0.12)
|
||||
self.th=threading.Thread(target=run,daemon=True); self.th.start()
|
||||
def stop(self):
|
||||
self.stop_evt.set()
|
||||
if self.th: self.th.join(timeout=0.4)
|
||||
if TTY: sys.stdout.write("\r\033[K"); sys.stdout.flush()
|
||||
|
||||
def stream_turn(p, sentinel, on_bytes):
|
||||
"""legge fino alla sentinella; on_bytes riceve i chunk della risposta. Poi legge la riga STAT."""
|
||||
pend=b""
|
||||
while True:
|
||||
b=p.stdout.read(1)
|
||||
if b==b"": return None
|
||||
pend+=b
|
||||
if pend.endswith(sentinel):
|
||||
rest=pend[:-len(sentinel)]
|
||||
if rest: on_bytes(rest)
|
||||
line=p.stdout.readline().decode("utf-8","replace").strip() # STAT tok tps hit rss
|
||||
m=re.match(r"STAT (\S+) (\S+) (\S+) (\S+)", line)
|
||||
return {"tok":int(m.group(1)),"tps":float(m.group(2)),"hit":float(m.group(3)),"rss":float(m.group(4))} if m else {}
|
||||
if len(pend)>len(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; <think></think> = risposta diretta (nothink)
|
||||
e=env_for(a); e["PROMPT"]=f"[gMASK]<sop><|user|>{prompt}<|assistant|><think></think>"
|
||||
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()
|
||||
@@ -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.<n_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()
|
||||
@@ -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()
|
||||
+145
@@ -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()
|
||||
@@ -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()
|
||||
@@ -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 <stdint.h>\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(cp<t[m][0]) hi=m-1; else if(cp>t[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""")
|
||||
+47
@@ -0,0 +1,47 @@
|
||||
/* Microbench: banda in lettura RANDOM con blocchi tipo-expert (~19 MB int4).
|
||||
* Misura cio' che fa il motore davvero: N thread che leggono in parallelo
|
||||
* (expert_load sotto omp parallel for), buffered oppure O_DIRECT.
|
||||
* uso: ./iobench <file_grande> [blocco_MB] [n_letture] [threads] [direct 0/1]
|
||||
* build: gcc -O2 -fopenmp iobench.c -o iobench */
|
||||
#define _GNU_SOURCE
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <stdint.h>
|
||||
#include <fcntl.h>
|
||||
#include <unistd.h>
|
||||
#include <time.h>
|
||||
#include <errno.h>
|
||||
#include <string.h>
|
||||
#include <omp.h>
|
||||
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<blk*2){fprintf(stderr,"file troppo piccolo\n");return 1;}
|
||||
/* offset random pre-generati (stessi per ogni configurazione: srand fisso) */
|
||||
off_t *offs=malloc(n*sizeof(off_t)); srand(1234);
|
||||
for(int i=0;i<n;i++){ off_t o=((off_t)rand()*4096)%(sz-blk); offs[i]=o&~4095L; }
|
||||
double t0=now(); long tot=0;
|
||||
#pragma omp parallel num_threads(nth) reduction(+:tot)
|
||||
{
|
||||
void *buf; if(posix_memalign(&buf,4096,blk)){perror("memalign");exit(1);}
|
||||
#pragma omp for schedule(dynamic,1)
|
||||
for(int i=0;i<n;i++){
|
||||
ssize_t r=pread(fd,buf,blk,offs[i]);
|
||||
if(r<0) perror("pread"); else tot+=r;
|
||||
}
|
||||
free(buf);
|
||||
}
|
||||
double dt=now()-t0;
|
||||
printf("%s x%d thread: %d letture x %ldMB = %.1f GB in %.2fs -> %.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;
|
||||
}
|
||||
@@ -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 <stdlib.h>
|
||||
#include <string.h>
|
||||
#include <stdio.h>
|
||||
#include <ctype.h>
|
||||
|
||||
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
|
||||
@@ -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")
|
||||
@@ -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 <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
#include <math.h>
|
||||
#include <time.h>
|
||||
#include <sys/resource.h>
|
||||
#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;kk<K;kk++) sm+=val[kk]; for(int kk=0;kk<K;kk++) val[kk]/=sm; }
|
||||
const float *xs = x + (int64_t)s*D;
|
||||
for (int kk = 0; kk < K; kk++) {
|
||||
Slot *e; expert_get(m, layer, idx[kk], &e);
|
||||
matmul_q(g, xs, e->g, e->gs, D, I); /* 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=<dir snapshot>\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<nfull;i++) printf("%d ", full[i]);
|
||||
printf("\nMotore C : "); for (int i=np;i<nfull;i++) { printf("%d ", out[i]); if (out[i]==full[i]) match++; }
|
||||
printf("\nToken coincidenti: %d/%d\n", match, n_new);
|
||||
double tot = m.hits + m.miss;
|
||||
printf("\nRSS PICCO: %.2f GB\n", rss_gb());
|
||||
printf("Hit-rate cache expert: %.1f%% (hit=%llu miss=%llu)\n", tot?100.0*m.hits/tot:0.0,
|
||||
(unsigned long long)m.hits, (unsigned long long)m.miss);
|
||||
printf("Velocita': %.2f tok/s (%.1fs per %d token)\n", n_new/dt, dt, n_new);
|
||||
free(buf); free(arena);
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1 @@
|
||||
{"prompt_ids": [3, 14, 159, 26, 53, 58, 200, 11, 77, 240, 5, 99], "full_ids": [3, 14, 159, 26, 53, 58, 200, 11, 77, 240, 5, 99, 207, 187, 119, 103, 103, 103, 103, 103, 119, 34, 103, 103, 103, 103, 103, 136, 112, 7, 119, 34], "tf_pred": [139, 197, 123, 34, 197, 34, 197, 197, 197, 193, 193, 207, 187, 119, 103, 103, 103, 103, 103, 119, 34, 103, 103, 103, 103, 103, 136, 112, 7, 119, 34, 103]}
|
||||
@@ -0,0 +1,38 @@
|
||||
#!/usr/bin/env bash
|
||||
# Pipeline GLM-5.2 (int4, streaming, 15 GB RAM) — tutto in WSL, modello su ext4.
|
||||
# uso: ./run.sh ["prompt"] [n_token]
|
||||
# Fa: (1) attende lo spostamento su ext4, (2) riprende la conversione fino a completarla,
|
||||
# (3) compila il motore, (4) genera testo restando nel budget RAM.
|
||||
set -euo pipefail
|
||||
|
||||
DIR=/home/vincenzo/glm52_i4 # modello int4 su ext4 (NON /mnt/c!)
|
||||
REPO=zai-org/GLM-5.2-FP8
|
||||
CODE=/mnt/c/Users/User/Desktop/moe-stream/c
|
||||
RAM_GB=15
|
||||
PROMPT="${1:-Ciao, chi sei?}"
|
||||
NGEN="${2:-64}"
|
||||
|
||||
cd "$CODE"
|
||||
|
||||
# 0) sanity: il modello deve stare su ext4, non su 9p/Windows
|
||||
case "$DIR" in /mnt/*) echo "ERRORE: $DIR e' su /mnt (9p/Windows). Mettilo su ext4."; exit 1;; esac
|
||||
|
||||
# 1) se un rsync di spostamento e' ancora vivo, aspettalo
|
||||
while pgrep -f "rsync.*glm52_i4" >/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
|
||||
+35
@@ -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 <dir-su-ext4/NVMe>
|
||||
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 <GB>"
|
||||
echo
|
||||
echo "IMPORTANTE: tieni il modello su disco VELOCE (NVMe/ext4), MAI su /mnt/c o rete."
|
||||
@@ -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 <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
#include <stdint.h>
|
||||
#include <fcntl.h>
|
||||
#include <unistd.h>
|
||||
#include <dirent.h>
|
||||
#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
|
||||
@@ -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
|
||||
@@ -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 <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
#include <stdint.h>
|
||||
#include <limits.h>
|
||||
#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;i<n;i++){ h^=(unsigned char)s[i]; h*=1099511628211ULL; } return h; }
|
||||
static void hm_init(hmap *m, int cap){ m->cap=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<len){ *cp=((c&0x1F)<<6)|(s[i+1]&0x3F); return 2; }
|
||||
if((c>>4)==0xE && i+2<len){ *cp=((c&0x0F)<<12)|((s[i+1]&0x3F)<<6)|(s[i+2]&0x3F); return 3; }
|
||||
if((c>>3)==0x1E && i+3<len){ *cp=((c&0x07)<<18)|((s[i+1]&0x3F)<<12)|((s[i+2]&0x3F)<<6)|(s[i+3]&0x3F); return 4; }
|
||||
*cp=c; return 1; /* byte invalido: trattato come singolo */
|
||||
}
|
||||
static int u8_put(char *o, uint32_t cp){
|
||||
if(cp<0x80){ o[0]=cp; return 1; }
|
||||
if(cp<0x800){ o[0]=0xC0|(cp>>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;i<vocab->len;i++){ int id=(int)vocab->kids[i]->num; if(id>maxid)maxid=id; }
|
||||
if(added) for(int i=0;i<added->len;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;i<vocab->len;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;i<merges->len;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;i<added->len;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;i<b;i++){ int bb=p[i]; memcpy(s+sl,T->byte2str[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(*no<max) out[(*no)++]=whole; free(s); return; }
|
||||
/* simboli iniziali = codepoint della stringa byte-level */
|
||||
int *soff=malloc((sl+1)*sizeof(int)), *slen=malloc((sl+1)*sizeof(int)); int ns=0;
|
||||
for(int i=0;i<sl;){ uint32_t cp; int k=u8_next((const unsigned char*)s,sl,i,&cp);
|
||||
soff[ns]=i; slen[ns]=k; ns++; i+=k; }
|
||||
char *kbuf=malloc(2*sl+2);
|
||||
for(;;){
|
||||
int best=INT_MAX, bp=-1;
|
||||
for(int i=0;i+1<ns;i++){
|
||||
int ll=slen[i], rl=slen[i+1];
|
||||
memcpy(kbuf,s+soff[i],ll); kbuf[ll]=0; memcpy(kbuf+ll+1,s+soff[i+1],rl);
|
||||
int rk=hm_get(&T->merges,kbuf,ll+1+rl);
|
||||
if(rk>=0 && rk<best){ best=rk; bp=i; }
|
||||
}
|
||||
if(bp<0) break;
|
||||
slen[bp]=soff[bp+1]+slen[bp+1]-soff[bp]; /* fonde bp e bp+1 (contigui in s) */
|
||||
for(int j=bp+1;j<ns-1;j++){ soff[j]=soff[j+1]; slen[j]=slen[j+1]; }
|
||||
ns--;
|
||||
}
|
||||
for(int i=0;i<ns;i++){
|
||||
int id=hm_get(&T->vocab,s+soff[i],slen[i]);
|
||||
if(id>=0 && *no<max) out[(*no)++]=id;
|
||||
}
|
||||
free(s); free(soff); free(slen); free(kbuf);
|
||||
}
|
||||
|
||||
/* ---------- pre-tokenizer regex (pattern cl100k) su una porzione di testo ----------
|
||||
* Decodifica i codepoint, applica le alternative IN ORDINE, e per ogni pezzo chiama bpe_piece. */
|
||||
static void pretok_chunk(Tok *T, const unsigned char *p, int a, int b, int *out, int *no, int max){
|
||||
int nb=b-a; if(nb<=0) return;
|
||||
uint32_t *cp=malloc((nb+1)*sizeof(uint32_t)); int *off=malloc((nb+2)*sizeof(int)); int n=0;
|
||||
for(int i=a;i<b;){ uint32_t c; int k=u8_next(p,b,i,&c); off[n]=i; cp[n]=c; n++; i+=k; }
|
||||
off[n]=b;
|
||||
#define ISNL(c) ((c)=='\r'||(c)=='\n')
|
||||
#define LOW(c) (((c)>='A'&&(c)<='Z')?((c)+32):(c))
|
||||
int i=0;
|
||||
while(i<n){
|
||||
int start=i; uint32_t c=cp[i];
|
||||
/* 1) (?i:'s|'t|'re|'ve|'m|'ll|'d) */
|
||||
if(c=='\'' && i+1<n){
|
||||
uint32_t d=LOW(cp[i+1]);
|
||||
if(i+2<n){ uint32_t d2=LOW(cp[i+2]);
|
||||
if((d=='r'&&d2=='e')||(d=='v'&&d2=='e')||(d=='l'&&d2=='l')){ i+=3; bpe_piece(T,p,off[start],off[i],out,no,max); continue; } }
|
||||
if(d=='s'||d=='t'||d=='m'||d=='d'){ i+=2; bpe_piece(T,p,off[start],off[i],out,no,max); continue; }
|
||||
}
|
||||
/* 2) [^\r\n\p{L}\p{N}]? \p{L}+ */
|
||||
{
|
||||
int j=i;
|
||||
if(!is_L(c) && !ISNL(c) && !is_N(c)){ if(j+1<n && is_L(cp[j+1])) j++; else j=-1; }
|
||||
if(j>=0){
|
||||
if(is_L(cp[j])){ while(j<n && is_L(cp[j])) j++; i=j; bpe_piece(T,p,off[start],off[i],out,no,max); continue; }
|
||||
}
|
||||
}
|
||||
/* 3) \p{N}{1,3} */
|
||||
if(is_N(c)){ int j=i,k=0; while(j<n && is_N(cp[j]) && k<3){ j++; k++; } i=j; bpe_piece(T,p,off[start],off[i],out,no,max); continue; }
|
||||
/* 4) ' ?[^\s\p{L}\p{N}]+[\r\n]*' */
|
||||
{
|
||||
int j=i;
|
||||
if(c==' ' && j+1<n && !is_S(cp[j+1]) && !is_L(cp[j+1]) && !is_N(cp[j+1])) j++;
|
||||
if(j<n && !is_S(cp[j]) && !is_L(cp[j]) && !is_N(cp[j])){
|
||||
while(j<n && !is_S(cp[j]) && !is_L(cp[j]) && !is_N(cp[j])) j++;
|
||||
while(j<n && ISNL(cp[j])) j++;
|
||||
i=j; bpe_piece(T,p,off[start],off[i],out,no,max); continue;
|
||||
}
|
||||
}
|
||||
/* 5) \s*[\r\n]+ -> run di whitespace fino all'ultimo newline contiguo */
|
||||
{
|
||||
int r=i; while(r<n && is_S(cp[r])) r++;
|
||||
if(r>i){ int last=-1; for(int j=i;j<r;j++) if(ISNL(cp[j])) last=j;
|
||||
if(last>=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<n) ? r-1 : r;
|
||||
if(end<=i) end=i+1; /* \s+ minimo 1 (fallback alt 7) */
|
||||
i=end; bpe_piece(T,p,off[start],off[i],out,no,max); continue;
|
||||
}
|
||||
}
|
||||
i++; /* salvagente: non dovrebbe accadere */
|
||||
bpe_piece(T,p,off[start],off[i],out,no,max);
|
||||
}
|
||||
#undef ISNL
|
||||
#undef LOW
|
||||
free(cp); free(off);
|
||||
}
|
||||
|
||||
/* ---------- encode: testo -> 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<len){
|
||||
/* prossima occorrenza di un added-token a partire da >= i (match piu' lungo) */
|
||||
int hitpos=-1, hitlen=0, hitid=-1;
|
||||
for(int j=i;j<len && hitpos<0;j++){
|
||||
for(int k=0;k<T->nsp;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<max) out[no++]=hitid;
|
||||
i=hitpos+hitlen;
|
||||
}
|
||||
return no;
|
||||
}
|
||||
|
||||
/* id di un added-token dato il suo contenuto (es. "<|endoftext|>"); -1 se assente */
|
||||
static int tok_id_of(Tok *T, const char *content){
|
||||
for(int i=0;i<T->nsp;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<n;i++){
|
||||
int id=ids[i]; if(id<0||id>=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;j<l && o<max;j++) out[o++]=s[j]; continue; }
|
||||
int sl=(int)strlen(s);
|
||||
for(int j=0;j<sl;){ uint32_t c; int k=u8_next((const unsigned char*)s,sl,j,&c); j+=k;
|
||||
if(c<1024 && T->cp2byte[c]>=0 && o<max) out[o++]=(char)(unsigned char)T->cp2byte[c]; }
|
||||
}
|
||||
if(o<max) out[o]=0;
|
||||
return o;
|
||||
}
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,46 @@
|
||||
/* Validazione del tokenizer C contro l'oracolo HF.
|
||||
* build: gcc -O2 tok_test.c -o tok_test
|
||||
* uso: ./tok_test <tokenizer.json> (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<ng&&ok;i++) ok = (got[i]==exp[i]);
|
||||
tot++; if(ok) pass++;
|
||||
/* round-trip decode */
|
||||
char dec[8192]; int dn=tok_decode(&T,got,ng,dec,8191);
|
||||
int drt = (dn==tn) && !memcmp(dec,tbuf,tn);
|
||||
if(drt) dpass++;
|
||||
if(!ok || !drt){
|
||||
fprintf(stderr,"MISMATCH text=%s\n exp(%d):",text,ne); for(int i=0;i<ne;i++)fprintf(stderr," %d",exp[i]);
|
||||
fprintf(stderr,"\n got(%d):",ng); for(int i=0;i<ng;i++)fprintf(stderr," %d",got[i]);
|
||||
fprintf(stderr,"\n decode_ok=%d\n", drt);
|
||||
}
|
||||
}
|
||||
printf("ENCODE: %d/%d DECODE(round-trip): %d/%d\n", pass,tot, dpass,tot);
|
||||
return pass==tot ? 0 : 2;
|
||||
}
|
||||
+162
@@ -0,0 +1,162 @@
|
||||
/* GENERATO da gen_unicode.py — non modificare a mano. */
|
||||
#ifndef TOK_UNICODE_H
|
||||
#define TOK_UNICODE_H
|
||||
#include <stdint.h>
|
||||
|
||||
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(cp<t[m][0]) hi=m-1; else if(cp>t[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
|
||||
Reference in New Issue
Block a user